AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

[rocprofv3] Add rocpd output support (part 1: prelude) (#401)

Просмотр исходного кода 
* [rocprofv3] Add rocpd output support (part 1: prelude)

- git submodules for sqlite3, GOTCHA, and pybind11
- HIP stream data
- rocprofiler_query_intercept_table_name(...)
- serialization load
- rocprofiler::sdk::get_perfetto_category(KindT)
- rocprofiler::sdk::parse::strip
- common library updates
  - md5sum
  - hasher
  - simple_timer
  - static_tl_object
  - get_process_start_time_ns(pid_t)
- output library updates
  - node_info
  - file_generator (generator is now virtual base class)
  - stream info updates

* Added submodules

* Code review updates

* Minor unused-but-set-X warning fixes

* Update CI

- install libsqlite3-dev package

* Update CI

- install libsqlite3-dev package

* Fix static thread-local object memory leak

- also fix signal handler chaining

* Remove URL from comment

* Remove page migration exception

* Enable ROCPROFILER_BUILD_SQLITE3 by default

- try find_package(SQLite3) first and then build when ROCPROFILER_BUILD_SQLITE3=ON

* Fix gotcha installation

- make install of target optional

* Validate tracing + counter collection dispatch data

- i.e. correlation ids, thread ids, timestamps

* Make find_package(SQLite3) optional

- ROCm CI does not have SQLite3 dev package installed and cannot build from source (missing tclsh)

* Fixes to tracing + counter collection test

* get_process_start_time_ns update

- original implementation did not work

* Fix pytest-packages test_perfetto_data for counter collection

- erroneous failure when used with same PMC + multiple agents

* cmake policy: option() honors normal variables

- for GOTCHA submodule

* Improve samples/api_buffered_tracing stability

- reduce likelihood of sporadic exception throw

* Update gotcha submodule

---------

Co-authored-by: Jonathan R. Madsen <jonathanrmadsen@gmail.com>

[ROCm/rocprofiler-sdk commit: 7166b1ab58]
Этот коммит содержится в:
Madsen, Jonathan
2025-05-18 20:11:26 -05:00
коммит произвёл GitHub
родитель 682b9967e0
Коммит b097e276a9
86 изменённых файлов: 4120 добавлений и 1730 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/rocprofiler-sdk/.github/workflows/codeql.yml
поставляемый
+2 -2
Просмотреть файл
@@ -62,11 +62,11 @@ jobs:
run: |
git config --global --add safe.directory '*'
apt-get update
apt-get install -y build-essential cmake g++-11 g++-12 python3-pip libdw-dev rccl-dev rccl-unittests
apt-get install -y build-essential cmake g++-11 g++-12 python3-pip libdw-dev libsqlite3-dev rccl-dev rccl-unittests
update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 10 --slave /usr/bin/g++ g++ /usr/bin/g++-11 --slave /usr/bin/gcov gcov /usr/bin/gcov-11
update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-12 20 --slave /usr/bin/g++ g++ /usr/bin/g++-12 --slave /usr/bin/gcov gcov /usr/bin/gcov-12
python3 -m pip install -r requirements.txt
# Initializes the CodeQL tools for scanning.
- name: Initialize CodeQL
projects/rocprofiler-sdk/.github/workflows/continuous_integration.yml
поставляемый
+4 -4
Просмотреть файл
@@ -65,7 +65,7 @@ jobs:
run: |
git config --global --add safe.directory '*'
apt-get update
apt-get install -y build-essential cmake g++-11 g++-12 python3-pip libdw-dev rccl-dev rccl-unittests rocjpeg-dev rocjpeg-test rocdecode-dev rocdecode-test
apt-get install -y build-essential cmake g++-11 g++-12 python3-pip libdw-dev libsqlite3-dev rccl-dev rccl-unittests rocjpeg-dev rocjpeg-test rocdecode-dev rocdecode-test
update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 10 --slave /usr/bin/g++ g++ /usr/bin/g++-11 --slave /usr/bin/gcov gcov /usr/bin/gcov-11
update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-12 20 --slave /usr/bin/g++ g++ /usr/bin/g++-12 --slave /usr/bin/gcov gcov /usr/bin/gcov-12
python3 -m pip install -U --user -r requirements.txt
@@ -348,7 +348,7 @@ jobs:
run: |
git config --global --add safe.directory '*'
apt-get update
apt-get install -y build-essential cmake python3-pip gcovr wkhtmltopdf xvfb xfonts-base xfonts-75dpi xfonts-100dpi xfonts-utils xfonts-encodings libfontconfig libdw-dev
apt-get install -y build-essential cmake python3-pip gcovr wkhtmltopdf xvfb xfonts-base xfonts-75dpi xfonts-100dpi xfonts-utils xfonts-encodings libfontconfig libdw-dev libsqlite3-dev
python3 -m pip install -U --user -r requirements.txt
rm -rf /opt/rocm/lib/*rocprofiler-sdk* /opt/rocm/lib/cmake/*rocprofiler-sdk* /opt/rocm/share/*rocprofiler-sdk* /opt/rocm/libexec/*rocprofiler-sdk*
@@ -612,11 +612,11 @@ jobs:
run: |
git config --global --add safe.directory '*'
apt-get update
apt-get install -y build-essential cmake python3-pip libasan8 libtsan2 software-properties-common clang-15
apt-get install -y build-essential cmake python3-pip libasan8 libtsan2 software-properties-common clang-15 libdw-dev libsqlite3-dev
add-apt-repository ppa:ubuntu-toolchain-r/test
apt-get update
apt-get upgrade -y
apt-get install -y gcc-${{ env.GCC_COMPILER_VERSION }} g++-${{ env.GCC_COMPILER_VERSION }} libdw-dev
apt-get install -y gcc-${{ env.GCC_COMPILER_VERSION }} g++-${{ env.GCC_COMPILER_VERSION }}
update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-${{ env.GCC_COMPILER_VERSION }} 100 --slave /usr/bin/g++ g++ /usr/bin/g++-${{ env.GCC_COMPILER_VERSION }} --slave /usr/bin/gcov gcov /usr/bin/gcov-${{ env.GCC_COMPILER_VERSION }}
python3 -m pip install -U --user -r requirements.txt
rm -rf /opt/rocm/lib/*rocprofiler-sdk* /opt/rocm/lib/cmake/*rocprofiler-sdk* /opt/rocm/share/*rocprofiler-sdk* /opt/rocm/libexec/*rocprofiler-sdk*
projects/rocprofiler-sdk/.github/workflows/docs.yml
поставляемый
+1 -1
Просмотреть файл
@@ -101,7 +101,7 @@ jobs:
run: |
git config --global --add safe.directory '*'
apt-get update
apt-get install -y cmake gcc g++ libdw-dev
apt-get install -y cmake gcc g++ libdw-dev libsqlite3-dev
python3 -m pip install -r requirements.txt
- name: Configure, Build, Install, and Package
projects/rocprofiler-sdk/.github/workflows/generate_afar.yml
поставляемый
+1 -1
Просмотреть файл
@@ -49,7 +49,7 @@ jobs:
apt-get update
amdgpu-repo --rocm-build=${{ inputs.build-job-name }}/${{ inputs.build-number }}
apt-get upgrade -y
apt-get install -y build-essential cmake rocm-llvm-dev python3-pip libdw-dev
apt-get install -y build-essential cmake rocm-llvm-dev python3-pip libdw-dev libsqlite3-dev
python3 -m pip install -r ${{github.workspace}}/rocprofiler-sdk/requirements.txt
- name: Building ROCProfiler-SDK Packages
projects/rocprofiler-sdk/.gitmodules
+9
Просмотреть файл
@@ -31,3 +31,12 @@
[submodule "external/json"]
path = external/json
url = https://github.com/nlohmann/json.git
[submodule "external/sqlite"]
path = external/sqlite
url = https://github.com/sqlite/sqlite
[submodule "external/pybind11"]
path = external/pybind11
url = https://github.com/pybind/pybind11.git
[submodule "external/gotcha"]
path = external/gotcha
url = https://jrmadsen@github.com/jrmadsen/GOTCHA
projects/rocprofiler-sdk/CMakeLists.txt
+4 -1
Просмотреть файл
@@ -80,7 +80,10 @@ set(CMAKE_BUILD_RPATH "${PROJECT_BINARY_DIR}/lib:\$ORIGIN:\$ORIGIN/../lib")
set(CMAKE_INSTALL_RPATH_USE_LINK_PATH
OFF
CACHE BOOL "")
mark_as_advanced(CMAKE_INSTALL_RPATH_USE_LINK_PATH)
set(CMAKE_DEBUG_POSTFIX
""
CACHE STRING "Suffix for non-executable build targets")
mark_as_advanced(CMAKE_INSTALL_RPATH_USE_LINK_PATH CMAKE_DEBUG_POSTFIX)
set(ROCPROFILER_INTERNAL_BUILD_DOCS
OFF
projects/rocprofiler-sdk/cmake/rocprofiler_build_settings.cmake
+13
Просмотреть файл
@@ -216,6 +216,19 @@ if(ROCPROFILER_BUILD_EXPERIMENTAL_WARNINGS)
INTERFACE rocprofiler-sdk::rocprofiler-sdk-experimental-flags)
endif()
# ----------------------------------------------------------------------------------------#
# extra flags for compiling with deprecated warnings
#
target_compile_definitions(rocprofiler-sdk-deprecated-flags
INTERFACE ROCPROFILER_SDK_DEPRECATED_WARNINGS=1)
rocprofiler_target_compile_options(rocprofiler-sdk-deprecated-flags
INTERFACE "-Wdeprecated-declarations")
if(ROCPROFILER_BUILD_DEPRECATED_WARNINGS)
target_link_libraries(rocprofiler-sdk-build-flags
INTERFACE rocprofiler-sdk::rocprofiler-sdk-deprecated-flags)
endif()
# ----------------------------------------------------------------------------------------#
# user customization
#
projects/rocprofiler-sdk/cmake/rocprofiler_interfaces.cmake
+9
Просмотреть файл
@@ -22,6 +22,12 @@ rocprofiler_add_interface_library(rocprofiler-sdk-perfetto "Enables Perfetto sup
rocprofiler_add_interface_library(rocprofiler-sdk-otf2 "Enables OTF2 support" INTERNAL)
rocprofiler_add_interface_library(rocprofiler-sdk-cereal "Enables Cereal support"
INTERNAL)
rocprofiler_add_interface_library(rocprofiler-sdk-sqlite3 "Enables SQLite3 support"
INTERNAL)
rocprofiler_add_interface_library(rocprofiler-sdk-pybind11 "Enables PyBind11 support"
INTERNAL)
rocprofiler_add_interface_library(rocprofiler-sdk-gotcha "Enables GOTCHA support"
INTERNAL)
rocprofiler_add_interface_library(rocprofiler-sdk-compile-definitions
"Compile definitions" INTERNAL)
rocprofiler_add_interface_library(rocprofiler-sdk-static-libgcc
@@ -41,6 +47,9 @@ rocprofiler_add_interface_library(rocprofiler-sdk-memcheck INTERFACE INTERNAL)
rocprofiler_add_interface_library(
rocprofiler-sdk-experimental-flags
"Compiler flags for experimental feature compilation" INTERNAL)
rocprofiler_add_interface_library(
rocprofiler-sdk-deprecated-flags
"Compiler flags for deprecated feature usage warnings" INTERNAL)
#
# interfaces for libraries (general)
projects/rocprofiler-sdk/cmake/rocprofiler_options.cmake
+8
Просмотреть файл
@@ -54,6 +54,12 @@ rocprofiler_add_option(
rocprofiler_add_option(ROCPROFILER_BUILD_FMT "Enable building fmt library internally" ON)
rocprofiler_add_option(ROCPROFILER_BUILD_GLOG
"Enable building glog (Google logging) library internally" ON)
rocprofiler_add_option(ROCPROFILER_BUILD_SQLITE3
"Enable building sqlite3 library internally" OFF)
rocprofiler_add_option(ROCPROFILER_BUILD_PYBIND11
"Enable building pybind11 library internally" ON)
rocprofiler_add_option(ROCPROFILER_BUILD_GOTCHA
"Enable building gotcha library internally" ON)
if(ROCPROFILER_BUILD_TESTS)
rocprofiler_add_option(
ROCPROFILER_BUILD_GTEST
@@ -87,6 +93,8 @@ rocprofiler_add_option(
"Enable warnings for experimental features but hide with -Wno-deprecated-declarations (this ensures that experimental warning message does not break macros)"
OFF
ADVANCED)
rocprofiler_add_option(ROCPROFILER_BUILD_DEPRECATED_WARNINGS
"Enable warnings for use of deprecated features" OFF ADVANCED)
# In the future, we will do this even with clang-tidy enabled
foreach(_OPT ROCPROFILER_BUILD_DEVELOPER ROCPROFILER_BUILD_WERROR)
projects/rocprofiler-sdk/docker/Dockerfile
+1 -1
Просмотреть файл
@@ -8,7 +8,7 @@ ARG BRANCH=amd-staging
RUN git clone -b ${BRANCH} https://github.com/ROCm/rocprofiler-sdk.git rocprofiler-sdk-source && \
python3 -m pip install -r rocprofiler-sdk-source/requirements.txt && \
sudo apt update && \
sudo apt install -y libdw-dev && \
sudo apt install -y libdw-dev libsqlite3-dev && \
cmake -B rocprofiler-sdk-build -DCMAKE_BUILD_TYPE=RelWithDebInfo -DROCPROFILER_BUILD_{SAMPLES,TESTS,BENCHMARK}=ON -DPython3_EXECUTABLE=$(which python3) -DCMAKE_INSTALL_PREFIX=$(realpath /opt/rocm) rocprofiler-sdk-source && \
cmake --build rocprofiler-sdk-build --target all --parallel 16 && \
sudo cmake --build rocprofiler-sdk-build --target install && \
projects/rocprofiler-sdk/external/CMakeLists.txt
поставляемый
+121 -2
Просмотреть файл
@@ -5,6 +5,9 @@
#
include(rocprofiler_utilities)
cmake_policy(PUSH)
cmake_policy(SET CMP0077 NEW) # option() honors normal variables
set(BUILD_TESTING OFF)
set(BUILD_SHARED_LIBS OFF)
set(BUILD_OBJECT_LIBS OFF) # Specific to PTL
@@ -207,7 +210,9 @@ if(TARGET rocprofiler-sdk-elfio)
INTERFACE ${ELFIO_INCLUDE_DIR})
endif()
#
# OTF2
#
add_subdirectory(otf2)
rocprofiler_checkout_git_submodule(
@@ -217,9 +222,123 @@ rocprofiler_checkout_git_submodule(
REPO_URL https://github.com/nlohmann/json.git
REPO_BRANCH "develop")
#
# JSON
#
add_subdirectory(json)
# target_link_libraries(rocprofiler-sdk-json INTERFACE $<BUILD_INTERFACE:json::json>)
target_include_directories(
rocprofiler-sdk-json SYSTEM
INTERFACE $<BUILD_INTERFACE:${PROJECT_SOURCE_DIR}/external/json/include>)
#
# SQLite3
#
if(ROCPROFILER_BUILD_SQLITE3)
# checkout submodule if not already checked out or clone repo if no .gitmodules file
rocprofiler_checkout_git_submodule(
RECURSIVE
RELATIVE_PATH external/sqlite
WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}
TEST_FILE configure
REPO_URL https://github.com/sqlite/sqlite
REPO_BRANCH "version-3.47.0")
find_program(
MAKE_COMMAND
NAMES make gmake
PATH_SUFFIXES bin REQUIRED)
include(ExternalProject)
externalproject_add(
rocprofiler-sdk-sqlite-build
PREFIX ${PROJECT_BINARY_DIR}/external/sqlite/build
SOURCE_DIR ${PROJECT_SOURCE_DIR}/external/sqlite
BUILD_IN_SOURCE 0
CONFIGURE_COMMAND
<SOURCE_DIR>/configure --prefix=${PROJECT_BINARY_DIR}/external/sqlite/install
--libdir=${PROJECT_BINARY_DIR}/external/sqlite/install/lib
--enable-tempstore=yes --enable-all --disable-tcl --with-pic CFLAGS=-O3\ -g1
BUILD_COMMAND ${MAKE_COMMAND} install -s
INSTALL_COMMAND "")
target_link_libraries(
rocprofiler-sdk-sqlite3
INTERFACE
$<BUILD_INTERFACE:${PROJECT_BINARY_DIR}/external/sqlite/install/lib/libsqlite3.so>
)
target_include_directories(
rocprofiler-sdk-sqlite3 SYSTEM
INTERFACE $<BUILD_INTERFACE:${PROJECT_BINARY_DIR}/external/sqlite/install/include>
)
add_dependencies(rocprofiler-sdk-sqlite3 rocprofiler-sdk-sqlite-build)
else()
# make it optional temporarily because ROCm CI does not have SQLite3 dev installed and
# cannot build from source (missing tclsh)
find_package(SQLite3)
if(SQLite3_FOUND)
target_link_libraries(rocprofiler-sdk-sqlite3 INTERFACE SQLite::SQLite3)
endif()
endif()
#
# PyBind11
#
if(ROCPROFILER_BUILD_PYBIND11)
# checkout submodule if not already checked out or clone repo if no .gitmodules file
rocprofiler_checkout_git_submodule(
RECURSIVE
RELATIVE_PATH external/pybind11
WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}
REPO_URL https://github.com/pybind/pybind11.git
REPO_BRANCH "v2.9.2")
target_include_directories(
rocprofiler-sdk-pybind11 SYSTEM
INTERFACE $<BUILD_INTERFACE:${PROJECT_SOURCE_DIR}/external/pybind11/include>)
else()
find_package(pybind11 REQUIRED)
target_link_libraries(rocprofiler-sdk-pybind11 INTERFACE pybind11::pybind11)
endif()
#
# GOTCHA
#
if(ROCPROFILER_BUILD_GOTCHA)
# checkout submodule if not already checked out or clone repo if no .gitmodules file
rocprofiler_checkout_git_submodule(
RECURSIVE
RELATIVE_PATH external/gotcha
WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}
REPO_URL https://github.com/jrmadsen/GOTCHA.git
REPO_BRANCH "rocprofiler")
function(add_gotcha_subdirectory)
set(GOTCHA_ENABLE_TESTS OFF)
set(GOTCHA_ENABLE_WARNING_ERROR OFF)
set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY
"${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_LIBDIR}/rocprofiler-sdk")
set(CMAKE_LIBRARY_OUTPUT_DIRECTORY
"${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_LIBDIR}/rocprofiler-sdk")
set(CMAKE_RUNTIME_OUTPUT_DIRECTORY
"${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_BINDIR}/rocprofiler-sdk")
add_subdirectory(gotcha EXCLUDE_FROM_ALL)
add_dependencies(rocprofiler-sdk-gotcha gotcha)
install(
TARGETS gotcha
DESTINATION "${CMAKE_INSTALL_LIBDIR}/rocprofiler-sdk"
OPTIONAL)
endfunction()
add_gotcha_subdirectory()
target_link_libraries(rocprofiler-sdk-gotcha INTERFACE gotcha::gotcha)
else()
find_package(gotcha REQUIRED)
if(TARGET gotcha::gotcha)
target_link_libraries(rocprofiler-sdk-gotcha INTERFACE gotcha::gotcha)
elseif(TARGET gotcha)
target_link_libraries(rocprofiler-sdk-gotcha INTERFACE gotcha)
endif()
endif()
cmake_policy(POP)
projects/rocprofiler-sdk/external/gotcha
поставляемый Подмодуль
+1
Submodule projects/rocprofiler-sdk/external/gotcha added at 9afbaf0ae5
projects/rocprofiler-sdk/external/pybind11
поставляемый Подмодуль
+1
Submodule projects/rocprofiler-sdk/external/pybind11 added at 914c06fb25
projects/rocprofiler-sdk/external/sqlite
поставляемый Подмодуль
+1
Submodule projects/rocprofiler-sdk/external/sqlite added at f5fb820c0f
projects/rocprofiler-sdk/samples/api_buffered_tracing/client.cpp
+28 -12
Просмотреть файл
@@ -56,6 +56,7 @@
#include <map>
#include <mutex>
#include <sstream>
#include <stdexcept>
#include <string>
#include <string_view>
#include <thread>
@@ -129,7 +130,8 @@ tool_code_object_callback(rocprofiler_callback_tracing_record_t record,
}
else if(record.phase == ROCPROFILER_CALLBACK_PHASE_UNLOAD)
{
client_kernels.erase(data->kernel_id);
// do not erase just in case a buffer callback needs this
// client_kernels.erase(data->kernel_id);
}
}
@@ -177,6 +179,21 @@ tool_tracing_callback(rocprofiler_context_id_t context,
}
}
auto get_name = [](const auto* _record) -> std::string_view {
try
{
return client_name_info.at(_record->kind, _record->operation);
} catch(std::exception& e)
{
std::cerr << __FUNCTION__
<< " threw an exception for buffer tracing kind=" << _record->kind
<< ", operation=" << _record->operation << "\nException: " << e.what()
<< std::flush;
abort();
}
return std::string_view{"??"};
};
if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING &&
(header->kind == ROCPROFILER_BUFFER_TRACING_HSA_CORE_API ||
header->kind == ROCPROFILER_BUFFER_TRACING_HSA_AMD_EXT_API ||
@@ -192,7 +209,7 @@ tool_tracing_callback(rocprofiler_context_id_t context,
<< ", extern_cid=" << record->correlation_id.external.value
<< ", kind=" << record->kind << ", operation=" << record->operation
<< ", start=" << record->start_timestamp << ", stop=" << record->end_timestamp
<< ", name=" << client_name_info.at(record->kind, record->operation);
<< ", name=" << get_name(record);
if(record->start_timestamp > record->end_timestamp)
{
@@ -242,6 +259,11 @@ tool_tracing_callback(rocprofiler_context_id_t context,
auto info = std::stringstream{};
auto kernel_id = record->dispatch_info.kernel_id;
auto kernel_name = (client_kernels.count(kernel_id) > 0)
? std::string_view{client_kernels.at(kernel_id).kernel_name}
: std::string_view{"??"};
info << "tid=" << record->thread_id << ", context=" << context.handle
<< ", buffer_id=" << buffer_id.handle
<< ", cid=" << record->correlation_id.internal
@@ -249,8 +271,7 @@ tool_tracing_callback(rocprofiler_context_id_t context,
<< ", kind=" << record->kind << ", operation=" << record->operation
<< ", agent_id=" << record->dispatch_info.agent_id.handle
<< ", queue_id=" << record->dispatch_info.queue_id.handle
<< ", kernel_id=" << record->dispatch_info.kernel_id
<< ", kernel=" << client_kernels.at(record->dispatch_info.kernel_id).kernel_name
<< ", kernel_id=" << record->dispatch_info.kernel_id << ", kernel=" << kernel_name
<< ", start=" << record->start_timestamp << ", stop=" << record->end_timestamp
<< ", private_segment_size=" << record->dispatch_info.private_segment_size
<< ", group_segment_size=" << record->dispatch_info.group_segment_size
@@ -282,8 +303,7 @@ tool_tracing_callback(rocprofiler_context_id_t context,
<< ", src_agent_id=" << record->src_agent_id.handle
<< ", dst_agent_id=" << record->dst_agent_id.handle
<< ", direction=" << record->operation << ", start=" << record->start_timestamp
<< ", stop=" << record->end_timestamp
<< ", name=" << client_name_info.at(record->kind, record->operation);
<< ", stop=" << record->end_timestamp << ", name=" << get_name(record);
if(record->start_timestamp > record->end_timestamp)
throw std::runtime_error("memory copy: start > end");
@@ -301,8 +321,7 @@ tool_tracing_callback(rocprofiler_context_id_t context,
info << "kind=" << record->kind << ", operation=" << record->operation
<< ", pid=" << record->pid << ", timestamp=" << record->timestamp
<< ", name=" << client_name_info.at(record->kind, record->operation)
<< std::boolalpha;
<< ", name=" << get_name(record) << std::boolalpha;
switch(record->operation)
{
@@ -382,8 +401,6 @@ tool_tracing_callback(rocprofiler_context_id_t context,
}
}
if(record->timestamp == 0) throw std::runtime_error("page migration: timestamp == 0");
static_cast<call_stack_t*>(user_data)->emplace_back(
source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()});
}
@@ -408,8 +425,7 @@ tool_tracing_callback(rocprofiler_context_id_t context,
<< ", agent_id=" << record->agent_id.handle
<< ", queue_id=" << record->queue_id.handle << ", thread_id=" << record->thread_id
<< ", elapsed=" << std::setprecision(3) << std::fixed << _elapsed
<< " usec, flags=" << record->flags
<< ", name=" << client_name_info.at(record->kind, record->operation);
<< " usec, flags=" << record->flags << ", name=" << get_name(record);
static_cast<call_stack_t*>(user_data)->emplace_back(
source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()});
projects/rocprofiler-sdk/samples/counter_collection/main.cpp
+5 -3
Просмотреть файл
@@ -106,6 +106,7 @@ launchKernels(const long NUM_LAUNCH, const long SYNC_INTERVAL, const int DEV_ID)
HIP_CALL(hipFree(gpuMem));
HIP_CALL(hipFree(A_d));
HIP_CALL(hipFree(C_d));
HIP_CALL(hipDeviceReset());
}
int
@@ -116,8 +117,8 @@ main(int argc, char** argv)
int ntotdevice = 0;
HIP_CALL(hipGetDeviceCount(&ntotdevice));
long nitr = 50000;
long nsync = 500;
long nitr = 5000;
long nsync = 50;
long ndevice = 0;
for(int i = 1; i < argc; ++i)
@@ -147,10 +148,11 @@ main(int argc, char** argv)
printf("[%s] Number of devices used: %li\n", exe_name, ndevice);
printf("[%s] Number of iterations: %li\n", exe_name, nitr);
printf("[%s] Syncing every %li iterations\n", exe_name, nsync);
std::cout << std::flush;
start();
for(long devid = 0; devid < ndevice; ++devid)
launchKernels(nitr, nsync, devid);
std::cerr << "Run complete\n";
std::cerr << "Run complete\n" << std::flush;
}
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/callback_tracing.h
+3 -2
Просмотреть файл
@@ -341,8 +341,9 @@ typedef struct rocprofiler_callback_tracing_runtime_initialization_data_t
*/
typedef struct rocprofiler_callback_tracing_hip_stream_data_t
{
uint64_t size; ///< size of this struct
rocprofiler_stream_id_t stream_id; ///< HIP stream ID
uint64_t size; ///< size of this struct
rocprofiler_stream_id_t stream_id; ///< HIP stream ID
rocprofiler_address_t stream_value; ///< HIP stream value
} rocprofiler_callback_tracing_hip_stream_data_t;
/**
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/CMakeLists.txt
+1
Просмотреть файл
@@ -13,3 +13,4 @@ install(
add_subdirectory(codeobj)
add_subdirectory(details)
add_subdirectory(serialization)
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/mpl.hpp
+6
Просмотреть файл
@@ -249,6 +249,12 @@ reserve(ContainerT& _c, ArgT _arg)
{
return impl::reserve(_c, 0, _arg);
}
template <typename Tp>
struct assert_false
{
static constexpr auto value = false;
};
} // namespace mpl
} // namespace sdk
} // namespace rocprofiler
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/tokenize.hpp
+19
Просмотреть файл
@@ -28,6 +28,7 @@
#include <sstream>
#include <stdexcept>
#include <string>
#include <string_view>
#include <vector>
namespace rocprofiler
@@ -211,6 +212,24 @@ str_transform(std::string_view input,
}
return _result;
}
inline std::string
strip(std::string&& str, std::string_view characters)
{
constexpr auto npos = std::string_view::npos;
auto bpos = str.find_first_not_of(characters);
auto epos = str.find_last_not_of(characters);
if(bpos != npos && epos != npos)
return str.substr(bpos, (epos - bpos + 1));
else if(bpos != npos)
return str.substr(bpos);
else if(epos != npos)
return str.substr(0, epos + 1);
return str;
}
} // namespace parse
} // namespace sdk
} // namespace rocprofiler
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp
+1
Просмотреть файл
@@ -115,6 +115,7 @@ get_enum_label(EnumT val, std::index_sequence<Idx, IdxTail...>)
return get_enum_label(val, std::index_sequence<IdxTail...>{});
return info::unsupported;
(void) val; // suppress unused-but-set-parameter warning
}
// clang-format off
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/perfetto.hpp
+113 -1
Просмотреть файл
@@ -65,6 +65,22 @@
.SetDescription( \
::rocprofiler::sdk::perfetto_category<::rocprofiler::sdk::TYPE>::description)
#define ROCPROFILER_PERFETTO_TRACING_CATEGORY(KIND, PREFIX, CODE, CATEGORY) \
template <> \
struct rocprofiler_tracing_perfetto_category<KIND, PREFIX##CODE> \
{ \
static constexpr auto name = \
::rocprofiler::sdk::perfetto_category<::rocprofiler::sdk::category::CATEGORY>::name; \
};
#define ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(...) \
ROCPROFILER_PERFETTO_TRACING_CATEGORY( \
rocprofiler_buffer_tracing_kind_t, ROCPROFILER_BUFFER_TRACING_, __VA_ARGS__)
#define ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(...) \
ROCPROFILER_PERFETTO_TRACING_CATEGORY( \
rocprofiler_callback_tracing_kind_t, ROCPROFILER_CALLBACK_TRACING_, __VA_ARGS__)
namespace rocprofiler
{
namespace sdk
@@ -85,6 +101,7 @@ ROCPROFILER_DEFINE_CATEGORY(category, memory_allocation, "Memory Allocation")
ROCPROFILER_DEFINE_CATEGORY(category, rocdecode_api, "rocDecode API function")
ROCPROFILER_DEFINE_CATEGORY(category, rocjpeg_api, "rocJPEG API function")
ROCPROFILER_DEFINE_CATEGORY(category, counter_collection, "Counter Collection")
ROCPROFILER_DEFINE_CATEGORY(category, none, "Unknown category")
#define ROCPROFILER_PERFETTO_CATEGORIES \
ROCPROFILER_PERFETTO_CATEGORY(category::hsa_api), \
@@ -97,7 +114,8 @@ ROCPROFILER_DEFINE_CATEGORY(category, counter_collection, "Counter Collection")
ROCPROFILER_PERFETTO_CATEGORY(category::counter_collection), \
ROCPROFILER_PERFETTO_CATEGORY(category::memory_allocation), \
ROCPROFILER_PERFETTO_CATEGORY(category::rocdecode_api), \
ROCPROFILER_PERFETTO_CATEGORY(category::rocjpeg_api)
ROCPROFILER_PERFETTO_CATEGORY(category::rocjpeg_api), \
ROCPROFILER_PERFETTO_CATEGORY(category::none)
#include <perfetto.h>
@@ -175,7 +193,101 @@ add_perfetto_annotation(perfetto_event_context_t& ctx, Np&& _name, Tp&& _val)
} // namespace sdk
} // namespace rocprofiler
#include <rocprofiler-sdk/fwd.h>
namespace rocprofiler
{
namespace sdk
{
template <typename KindT, size_t Idx>
struct rocprofiler_tracing_perfetto_category;
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(NONE, none)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(HSA_CORE_API, hsa_api)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(HSA_AMD_EXT_API, hsa_api)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(HSA_IMAGE_EXT_API, hsa_api)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(HSA_FINALIZE_EXT_API, hsa_api)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(HIP_RUNTIME_API, hip_api)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(HIP_COMPILER_API, hip_api)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(MARKER_CORE_API, marker_api)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(MARKER_CONTROL_API, marker_api)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(MARKER_NAME_API, marker_api)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(MEMORY_COPY, memory_copy)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(MEMORY_ALLOCATION, memory_allocation)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(KERNEL_DISPATCH, kernel_dispatch)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(PAGE_MIGRATION, none)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(SCRATCH_MEMORY, memory_allocation)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(CORRELATION_ID_RETIREMENT, none)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(RCCL_API, rccl_api)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(OMPT, openmp)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(RUNTIME_INITIALIZATION, none)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(ROCDECODE_API, rocdecode_api)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(ROCJPEG_API, rocjpeg_api)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(HIP_STREAM, hip_api)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(HIP_RUNTIME_API_EXT, hip_api)
ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY(HIP_COMPILER_API_EXT, hip_api)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(NONE, none)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(HSA_CORE_API, hsa_api)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(HSA_AMD_EXT_API, hsa_api)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(HSA_IMAGE_EXT_API, hsa_api)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(HSA_FINALIZE_EXT_API, hsa_api)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(HIP_RUNTIME_API, hip_api)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(HIP_COMPILER_API, hip_api)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(MARKER_CORE_API, marker_api)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(MARKER_CONTROL_API, marker_api)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(MARKER_NAME_API, marker_api)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(CODE_OBJECT, none)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(SCRATCH_MEMORY, memory_allocation)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(KERNEL_DISPATCH, kernel_dispatch)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(MEMORY_COPY, memory_copy)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(MEMORY_ALLOCATION, memory_allocation)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(RCCL_API, rccl_api)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(OMPT, openmp)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(RUNTIME_INITIALIZATION, none)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(ROCDECODE_API, rocdecode_api)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(ROCJPEG_API, rocjpeg_api)
ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY(HIP_STREAM, hip_api)
template <typename KindT, size_t Idx, size_t... Tail>
const char*
get_perfetto_category(KindT kind, std::index_sequence<Idx, Tail...>)
{
if(kind == Idx) return rocprofiler_tracing_perfetto_category<KindT, Idx>::name;
// recursion until tail empty
if constexpr(sizeof...(Tail) > 0)
return get_perfetto_category(kind, std::index_sequence<Tail...>{});
return nullptr;
}
template <typename KindT>
const char*
get_perfetto_category(KindT kind)
{
if constexpr(std::is_same<KindT, rocprofiler_buffer_tracing_kind_t>::value)
{
return get_perfetto_category<KindT>(
kind, std::make_index_sequence<ROCPROFILER_BUFFER_TRACING_LAST>{});
}
else if constexpr(std::is_same<KindT, rocprofiler_callback_tracing_kind_t>::value)
{
return get_perfetto_category<KindT>(
kind, std::make_index_sequence<ROCPROFILER_CALLBACK_TRACING_LAST>{});
}
else
{
static_assert(rocprofiler::sdk::mpl::assert_false<KindT>::value, "Unsupported data type");
}
return nullptr;
}
} // namespace sdk
} // namespace rocprofiler
#undef ROCPROFILER_DEFINE_PERFETTO_CATEGORY
#undef ROCPROFILER_DEFINE_CATEGORY
#undef ROCPROFILER_PERFETTO_CATEGORY
#undef ROCPROFILER_PERFETTO_CATEGORIES
#undef ROCPROFILER_PERFETTO_TRACING_CATEGORY
#undef ROCPROFILER_PERFETTO_BUFFER_TRACING_CATEGORY
#undef ROCPROFILER_PERFETTO_CALLBACK_TRACING_CATEGORY
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/serialization.hpp
+2 -1317
Просмотреть файл
Разница между файлами не показана из-за своего большого размера Загрузить разницу
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/serialization/CMakeLists.txt
+11
Просмотреть файл
@@ -0,0 +1,11 @@
#
#
# Installation of public C++ headers (implementations)
#
#
set(ROCPROFILER_CXX_SERIALIZATION_HEADER_FILES save.hpp load.hpp)
install(
FILES ${ROCPROFILER_CXX_SERIALIZATION_HEADER_FILES}
DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/rocprofiler-sdk/cxx/serialization
COMPONENT development)
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/serialization/load.hpp
+380
Просмотреть файл
@@ -0,0 +1,380 @@
// MIT License
//
// Copyright (c) 2023-2025 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/buffer.h>
#include <rocprofiler-sdk/callback_tracing.h>
#include <rocprofiler-sdk/external_correlation.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/internal_threading.h>
#include <rocprofiler-sdk/rocprofiler.h>
#include <rocprofiler-sdk/cxx/name_info.hpp>
#include <rocprofiler-sdk/cxx/utility.hpp>
#include <cereal/archives/binary.hpp>
#include <cereal/archives/json.hpp>
#include <cereal/archives/portable_binary.hpp>
#include <cereal/cereal.hpp>
#include <cereal/types/array.hpp>
#include <cereal/types/atomic.hpp>
#include <cereal/types/bitset.hpp>
#include <cereal/types/chrono.hpp>
#include <cereal/types/common.hpp>
#include <cereal/types/complex.hpp>
#include <cereal/types/deque.hpp>
#include <cereal/types/functional.hpp>
#include <cereal/types/list.hpp>
#include <cereal/types/map.hpp>
#include <cereal/types/optional.hpp>
#include <cereal/types/queue.hpp>
#include <cereal/types/set.hpp>
#include <cereal/types/stack.hpp>
#include <cereal/types/string.hpp>
#include <cereal/types/unordered_map.hpp>
#include <cereal/types/unordered_set.hpp>
#include <cereal/types/utility.hpp>
#include <cereal/types/variant.hpp>
#include <cereal/types/vector.hpp>
#include <string>
#include <string_view>
#include <type_traits>
#include <utility>
#include <vector>
#if defined(ROCPROFILER_SDK_CXX_SERIALIZATION_LOAD_DEBUG) && \
ROCPROFILER_SDK_CXX_SERIALIZATION_LOAD_DEBUG > 0
# define ROCP_SDK_LOAD_MESSAGE(NAME) \
std::clog << "[" << __PRETTY_FUNCTION__ << "] loading JSON field " << NAME << "...\n" \
<< std::flush
#else
# define ROCP_SDK_LOAD_MESSAGE(NAME)
#endif
#define ROCP_SDK_LOAD_DATA_FIELD(FIELD) \
ROCP_SDK_LOAD_MESSAGE(#FIELD); \
ar(make_nvp(#FIELD, data.FIELD))
#define ROCP_SDK_LOAD_DATA_VALUE(NAME, VALUE) \
ROCP_SDK_LOAD_MESSAGE(NAME); \
ar(make_nvp(NAME, data.VALUE))
#define ROCP_SDK_LOAD_VALUE(NAME, VALUE) \
ROCP_SDK_LOAD_MESSAGE(NAME); \
ar(make_nvp(NAME, VALUE))
#define ROCP_SDK_LOAD_DATA_CSTR(FIELD) \
{ \
ROCP_SDK_LOAD_MESSAGE(#FIELD); \
auto _val = new std::string{}; \
ar(make_nvp(#FIELD, *_val)); \
data.FIELD = _val->c_str(); \
}
#define ROCP_SDK_LOAD_DATA_BITFIELD(NAME, VALUE) \
{ \
ROCP_SDK_LOAD_MESSAGE(NAME); \
auto _val = data.VALUE; \
ar(make_nvp(NAME, _val)); \
data.VALUE = _val; \
}
#if !defined(ROCPROFILER_SDK_CEREAL_NAMESPACE_BEGIN)
# define ROCPROFILER_SDK_CEREAL_NAMESPACE_BEGIN \
namespace cereal \
{
#endif
#if !defined(ROCPROFILER_SDK_CEREAL_NAMESPACE_END)
# define ROCPROFILER_SDK_CEREAL_NAMESPACE_END } // namespace cereal
#endif
ROCPROFILER_SDK_CEREAL_NAMESPACE_BEGIN
template <typename ArchiveT>
void
load(ArchiveT& ar, rocprofiler_agent_id_t& data)
{
ROCP_SDK_LOAD_DATA_FIELD(handle);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, HsaCacheType& data)
{
ROCP_SDK_LOAD_DATA_BITFIELD("Data", ui32.Data);
ROCP_SDK_LOAD_DATA_BITFIELD("Instruction", ui32.Instruction);
ROCP_SDK_LOAD_DATA_BITFIELD("CPU", ui32.CPU);
ROCP_SDK_LOAD_DATA_BITFIELD("HSACU", ui32.HSACU);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, HSA_LINKPROPERTY& data)
{
ROCP_SDK_LOAD_DATA_BITFIELD("Override", ui32.Override);
ROCP_SDK_LOAD_DATA_BITFIELD("NonCoherent", ui32.NonCoherent);
ROCP_SDK_LOAD_DATA_BITFIELD("NoAtomics32bit", ui32.NoAtomics32bit);
ROCP_SDK_LOAD_DATA_BITFIELD("NoAtomics64bit", ui32.NoAtomics64bit);
ROCP_SDK_LOAD_DATA_BITFIELD("NoPeerToPeerDMA", ui32.NoPeerToPeerDMA);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, HSA_CAPABILITY& data)
{
ROCP_SDK_LOAD_DATA_BITFIELD("HotPluggable", ui32.HotPluggable);
ROCP_SDK_LOAD_DATA_BITFIELD("HSAMMUPresent", ui32.HSAMMUPresent);
ROCP_SDK_LOAD_DATA_BITFIELD("SharedWithGraphics", ui32.SharedWithGraphics);
ROCP_SDK_LOAD_DATA_BITFIELD("QueueSizePowerOfTwo", ui32.QueueSizePowerOfTwo);
ROCP_SDK_LOAD_DATA_BITFIELD("QueueSize32bit", ui32.QueueSize32bit);
ROCP_SDK_LOAD_DATA_BITFIELD("QueueIdleEvent", ui32.QueueIdleEvent);
ROCP_SDK_LOAD_DATA_BITFIELD("VALimit", ui32.VALimit);
ROCP_SDK_LOAD_DATA_BITFIELD("WatchPointsSupported", ui32.WatchPointsSupported);
ROCP_SDK_LOAD_DATA_BITFIELD("WatchPointsTotalBits", ui32.WatchPointsTotalBits);
ROCP_SDK_LOAD_DATA_BITFIELD("DoorbellType", ui32.DoorbellType);
ROCP_SDK_LOAD_DATA_BITFIELD("AQLQueueDoubleMap", ui32.AQLQueueDoubleMap);
ROCP_SDK_LOAD_DATA_BITFIELD("DebugTrapSupported", ui32.DebugTrapSupported);
ROCP_SDK_LOAD_DATA_BITFIELD("WaveLaunchTrapOverrideSupported",
ui32.WaveLaunchTrapOverrideSupported);
ROCP_SDK_LOAD_DATA_BITFIELD("WaveLaunchModeSupported", ui32.WaveLaunchModeSupported);
ROCP_SDK_LOAD_DATA_BITFIELD("PreciseMemoryOperationsSupported",
ui32.PreciseMemoryOperationsSupported);
ROCP_SDK_LOAD_DATA_BITFIELD("DEPRECATED_SRAM_EDCSupport", ui32.DEPRECATED_SRAM_EDCSupport);
ROCP_SDK_LOAD_DATA_BITFIELD("Mem_EDCSupport", ui32.Mem_EDCSupport);
ROCP_SDK_LOAD_DATA_BITFIELD("RASEventNotify", ui32.RASEventNotify);
ROCP_SDK_LOAD_DATA_BITFIELD("ASICRevision", ui32.ASICRevision);
ROCP_SDK_LOAD_DATA_BITFIELD("SRAM_EDCSupport", ui32.SRAM_EDCSupport);
ROCP_SDK_LOAD_DATA_BITFIELD("SVMAPISupported", ui32.SVMAPISupported);
ROCP_SDK_LOAD_DATA_BITFIELD("CoherentHostAccess", ui32.CoherentHostAccess);
ROCP_SDK_LOAD_DATA_BITFIELD("DebugSupportedFirmware", ui32.DebugSupportedFirmware);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, rocprofiler_dim3_t& data)
{
ROCP_SDK_LOAD_DATA_FIELD(x);
ROCP_SDK_LOAD_DATA_FIELD(y);
ROCP_SDK_LOAD_DATA_FIELD(z);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, HSA_MEMORYPROPERTY& data)
{
ROCP_SDK_LOAD_DATA_BITFIELD("HotPluggable", ui32.HotPluggable);
ROCP_SDK_LOAD_DATA_BITFIELD("NonVolatile", ui32.NonVolatile);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, HSA_ENGINE_VERSION& data)
{
ROCP_SDK_LOAD_DATA_BITFIELD("uCodeSDMA", uCodeSDMA);
ROCP_SDK_LOAD_DATA_BITFIELD("uCodeRes", uCodeRes);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, HSA_ENGINE_ID& data)
{
ROCP_SDK_LOAD_DATA_BITFIELD("uCode", ui32.uCode);
ROCP_SDK_LOAD_DATA_BITFIELD("Major", ui32.Major);
ROCP_SDK_LOAD_DATA_BITFIELD("Minor", ui32.Minor);
ROCP_SDK_LOAD_DATA_BITFIELD("Stepping", ui32.Stepping);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, rocprofiler_agent_cache_t& data)
{
ROCP_SDK_LOAD_DATA_FIELD(processor_id_low);
ROCP_SDK_LOAD_DATA_FIELD(size);
ROCP_SDK_LOAD_DATA_FIELD(level);
ROCP_SDK_LOAD_DATA_FIELD(cache_line_size);
ROCP_SDK_LOAD_DATA_FIELD(cache_lines_per_tag);
ROCP_SDK_LOAD_DATA_FIELD(association);
ROCP_SDK_LOAD_DATA_FIELD(latency);
ROCP_SDK_LOAD_DATA_FIELD(type);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, rocprofiler_agent_io_link_t& data)
{
ROCP_SDK_LOAD_DATA_FIELD(type);
ROCP_SDK_LOAD_DATA_FIELD(version_major);
ROCP_SDK_LOAD_DATA_FIELD(version_minor);
ROCP_SDK_LOAD_DATA_FIELD(node_from);
ROCP_SDK_LOAD_DATA_FIELD(node_to);
ROCP_SDK_LOAD_DATA_FIELD(weight);
ROCP_SDK_LOAD_DATA_FIELD(min_latency);
ROCP_SDK_LOAD_DATA_FIELD(max_latency);
ROCP_SDK_LOAD_DATA_FIELD(min_bandwidth);
ROCP_SDK_LOAD_DATA_FIELD(max_bandwidth);
ROCP_SDK_LOAD_DATA_FIELD(recommended_transfer_size);
ROCP_SDK_LOAD_DATA_FIELD(flags);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, rocprofiler_agent_mem_bank_t& data)
{
ROCP_SDK_LOAD_DATA_FIELD(heap_type);
ROCP_SDK_LOAD_DATA_FIELD(flags);
ROCP_SDK_LOAD_DATA_FIELD(width);
ROCP_SDK_LOAD_DATA_FIELD(mem_clk_max);
ROCP_SDK_LOAD_DATA_FIELD(size_in_bytes);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, rocprofiler_agent_v0_t& data)
{
ROCP_SDK_LOAD_DATA_FIELD(size);
ROCP_SDK_LOAD_DATA_FIELD(id);
ROCP_SDK_LOAD_DATA_FIELD(type);
ROCP_SDK_LOAD_DATA_FIELD(cpu_cores_count);
ROCP_SDK_LOAD_DATA_FIELD(simd_count);
ROCP_SDK_LOAD_DATA_FIELD(mem_banks_count);
ROCP_SDK_LOAD_DATA_FIELD(caches_count);
ROCP_SDK_LOAD_DATA_FIELD(io_links_count);
ROCP_SDK_LOAD_DATA_FIELD(cpu_core_id_base);
ROCP_SDK_LOAD_DATA_FIELD(simd_id_base);
ROCP_SDK_LOAD_DATA_FIELD(max_waves_per_simd);
ROCP_SDK_LOAD_DATA_FIELD(lds_size_in_kb);
ROCP_SDK_LOAD_DATA_FIELD(gds_size_in_kb);
ROCP_SDK_LOAD_DATA_FIELD(num_gws);
ROCP_SDK_LOAD_DATA_FIELD(wave_front_size);
ROCP_SDK_LOAD_DATA_FIELD(num_xcc);
ROCP_SDK_LOAD_DATA_FIELD(cu_count);
ROCP_SDK_LOAD_DATA_FIELD(array_count);
ROCP_SDK_LOAD_DATA_FIELD(num_shader_banks);
ROCP_SDK_LOAD_DATA_FIELD(simd_arrays_per_engine);
ROCP_SDK_LOAD_DATA_FIELD(cu_per_simd_array);
ROCP_SDK_LOAD_DATA_FIELD(simd_per_cu);
ROCP_SDK_LOAD_DATA_FIELD(max_slots_scratch_cu);
ROCP_SDK_LOAD_DATA_FIELD(gfx_target_version);
ROCP_SDK_LOAD_DATA_FIELD(vendor_id);
ROCP_SDK_LOAD_DATA_FIELD(device_id);
ROCP_SDK_LOAD_DATA_FIELD(location_id);
ROCP_SDK_LOAD_DATA_FIELD(domain);
ROCP_SDK_LOAD_DATA_FIELD(drm_render_minor);
ROCP_SDK_LOAD_DATA_FIELD(num_sdma_engines);
ROCP_SDK_LOAD_DATA_FIELD(num_sdma_xgmi_engines);
ROCP_SDK_LOAD_DATA_FIELD(num_sdma_queues_per_engine);
ROCP_SDK_LOAD_DATA_FIELD(num_cp_queues);
ROCP_SDK_LOAD_DATA_FIELD(max_engine_clk_ccompute);
ROCP_SDK_LOAD_DATA_FIELD(max_engine_clk_fcompute);
ROCP_SDK_LOAD_DATA_FIELD(sdma_fw_version);
ROCP_SDK_LOAD_DATA_FIELD(fw_version);
ROCP_SDK_LOAD_DATA_FIELD(capability);
ROCP_SDK_LOAD_DATA_FIELD(cu_per_engine);
ROCP_SDK_LOAD_DATA_FIELD(max_waves_per_cu);
ROCP_SDK_LOAD_DATA_FIELD(family_id);
ROCP_SDK_LOAD_DATA_FIELD(workgroup_max_size);
ROCP_SDK_LOAD_DATA_FIELD(grid_max_size);
ROCP_SDK_LOAD_DATA_FIELD(local_mem_size);
ROCP_SDK_LOAD_DATA_FIELD(hive_id);
ROCP_SDK_LOAD_DATA_FIELD(gpu_id);
ROCP_SDK_LOAD_DATA_FIELD(workgroup_max_dim);
ROCP_SDK_LOAD_DATA_FIELD(grid_max_dim);
ROCP_SDK_LOAD_DATA_CSTR(name);
ROCP_SDK_LOAD_DATA_CSTR(vendor_name);
ROCP_SDK_LOAD_DATA_CSTR(product_name);
ROCP_SDK_LOAD_DATA_CSTR(model_name);
ROCP_SDK_LOAD_DATA_FIELD(node_id);
ROCP_SDK_LOAD_DATA_FIELD(logical_node_id);
ROCP_SDK_LOAD_DATA_FIELD(logical_node_type_id);
auto generate = [&](auto name, const auto*& value, auto& size) {
using value_type =
std::remove_const_t<std::remove_pointer_t<std::remove_reference_t<decltype(value)>>>;
auto vec = std::vector<value_type>{};
ar(make_nvp(name, vec));
size = vec.size();
auto* value_m = new value_type[size];
for(uint64_t i = 0; i < size; ++i)
value_m[i] = vec.at(i);
value = value_m;
};
generate("mem_banks", data.mem_banks, data.mem_banks_count);
generate("caches", data.caches, data.caches_count);
generate("io_links", data.io_links, data.io_links_count);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, rocprofiler_address_t& data)
{
ROCP_SDK_LOAD_DATA_FIELD(handle);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, rocprofiler_callback_tracing_code_object_load_data_t& data)
{
ROCP_SDK_LOAD_DATA_FIELD(size);
ROCP_SDK_LOAD_DATA_FIELD(code_object_id);
ROCP_SDK_LOAD_DATA_FIELD(agent_id);
ROCP_SDK_LOAD_DATA_CSTR(uri);
ROCP_SDK_LOAD_DATA_FIELD(load_base);
ROCP_SDK_LOAD_DATA_FIELD(load_size);
ROCP_SDK_LOAD_DATA_FIELD(load_delta);
ROCP_SDK_LOAD_DATA_FIELD(storage_type);
if(data.storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_FILE)
{
ROCP_SDK_LOAD_DATA_FIELD(storage_file);
}
else if(data.storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_MEMORY)
{
ROCP_SDK_LOAD_DATA_FIELD(memory_base);
ROCP_SDK_LOAD_DATA_FIELD(memory_size);
}
}
template <typename ArchiveT>
void
load(ArchiveT& ar, rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t& data)
{
ROCP_SDK_LOAD_DATA_FIELD(size);
ROCP_SDK_LOAD_DATA_FIELD(kernel_id);
ROCP_SDK_LOAD_DATA_FIELD(code_object_id);
ROCP_SDK_LOAD_DATA_CSTR(kernel_name);
ROCP_SDK_LOAD_DATA_FIELD(kernel_object);
ROCP_SDK_LOAD_DATA_FIELD(kernarg_segment_size);
ROCP_SDK_LOAD_DATA_FIELD(kernarg_segment_alignment);
ROCP_SDK_LOAD_DATA_FIELD(group_segment_size);
ROCP_SDK_LOAD_DATA_FIELD(private_segment_size);
ROCP_SDK_LOAD_DATA_FIELD(sgpr_count);
ROCP_SDK_LOAD_DATA_FIELD(arch_vgpr_count);
ROCP_SDK_LOAD_DATA_FIELD(accum_vgpr_count);
ROCP_SDK_LOAD_DATA_FIELD(kernel_code_entry_byte_offset);
ROCP_SDK_LOAD_DATA_FIELD(kernel_address);
}
ROCPROFILER_SDK_CEREAL_NAMESPACE_END
#undef ROCP_SDK_LOAD_DATA_FIELD
#undef ROCP_SDK_LOAD_DATA_VALUE
#undef ROCP_SDK_LOAD_DATA_CSTR
#undef ROCP_SDK_LOAD_DATA_BITFIELD
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/serialization/save.hpp
+1342
Просмотреть файл
Разница между файлами не показана из-за своего большого размера Загрузить разницу
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/intercept_table.h
+22
Просмотреть файл
@@ -68,6 +68,28 @@ typedef void (*rocprofiler_intercept_library_cb_t)(rocprofiler_intercept_table_t
uint64_t num_tables,
void* user_data);
/**
* @brief (experimental) Query the name of the intercept table. The name retrieved from this
* function is a string literal that is encoded in the read-only section of the binary (i.e. it is
* always "allocated" and never "deallocated").
*
* @param [in] kind Intercept table kind
* @param [out] name If non-null and the name is a constant string that does not require dynamic
* allocation, this paramter will be set to the address of the string literal, otherwise it will
* be set to nullptr
* @param [out] name_len If non-null, this will be assigned the length of the name (regardless of
* the name is a constant string or requires dynamic allocation)
* @return ::rocprofiler_status_t
* @retval ::ROCPROFILER_STATUS_ERROR_KIND_NOT_FOUND Returned if the domain id is not valid
* @retval ::ROCPROFILER_STATUS_SUCCESS Returned if a valid domain, regardless if there is a
* constant string or not.
*/
ROCPROFILER_SDK_EXPERIMENTAL
rocprofiler_status_t
rocprofiler_query_intercept_table_name(rocprofiler_intercept_table_t kind,
const char** name,
uint64_t* name_len) ROCPROFILER_API;
/**
* @brief (experimental) Invoke this function to receive callbacks when a ROCm library registers its
* API intercept table with rocprofiler. Use the ::rocprofiler_intercept_table_t enumeration for
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/ompt/omp-tools.h
+5 -7
Просмотреть файл
@@ -104,10 +104,9 @@
macro(ompt_callback_thread_begin, ompt_callback_thread_begin_t, 1) /* thread begin */ \
macro(ompt_callback_thread_end, ompt_callback_thread_end_t, 2) /* thread end */ \
\
macro(ompt_callback_parallel_begin, \
ompt_callback_parallel_begin_t, \
3) /* parallel begin */ \
macro(ompt_callback_parallel_end, ompt_callback_parallel_end_t, 4) /* parallel end */ \
macro( \
ompt_callback_parallel_begin, ompt_callback_parallel_begin_t, 3) /* parallel begin */ \
macro(ompt_callback_parallel_end, ompt_callback_parallel_end_t, 4) /* parallel end */ \
\
macro(ompt_callback_task_create, ompt_callback_task_create_t, 5) /* task begin */ \
macro(ompt_callback_task_schedule, ompt_callback_task_schedule_t, 6) /* task schedule */ \
@@ -171,9 +170,8 @@
#define FOREACH_OMPT_NOEMI_EVENT(macro) \
/*--- Mandatory Events ---*/ \
macro(ompt_callback_target, ompt_callback_target_t, 8) /* target */ \
macro(ompt_callback_target_data_op, \
ompt_callback_target_data_op_t, \
9) /* target data op */ \
macro( \
ompt_callback_target_data_op, ompt_callback_target_data_op_t, 9) /* target data op */ \
macro(ompt_callback_target_submit, \
ompt_callback_target_submit_t, \
10) /* target submit */ /* Optional Events */ \
projects/rocprofiler-sdk/source/lib/common/CMakeLists.txt
+14 -2
Просмотреть файл
@@ -3,8 +3,17 @@
#
rocprofiler_activate_clang_tidy()
set(common_sources demangle.cpp elf_utils.cpp environment.cpp logging.cpp
static_object.cpp string_entry.cpp utility.cpp)
set(common_sources
demangle.cpp
elf_utils.cpp
environment.cpp
logging.cpp
md5sum.cpp
simple_timer.cpp
static_object.cpp
static_tl_object.cpp
string_entry.cpp
utility.cpp)
set(common_headers
abi.hpp
defines.hpp
@@ -12,9 +21,12 @@ set(common_headers
elf_utils.hpp
environment.hpp
filesystem.hpp
hasher.hpp
logging.hpp
md5sum.hpp
mpl.hpp
scope_destructor.hpp
simple_timer.hpp
static_object.hpp
static_tl_object.hpp
string_entry.hpp
projects/rocprofiler-sdk/source/lib/common/hasher.hpp
+135
Просмотреть файл
@@ -0,0 +1,135 @@
// MIT License
//
// Copyright (c) 2023-2025 Advanced Micro Devices, Inc.
//
// 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 "lib/common/defines.hpp"
#include <rocprofiler-sdk/cxx/details/mpl.hpp>
#include <cstddef>
#include <cstdint>
#include <string>
#include <type_traits>
namespace rocprofiler
{
namespace common
{
// A helper class which computes a 64-bit hash of the input data.
// The algorithm used is FNV-1a as it is fast and easy to implement and has
// relatively few collisions.
// WARNING: This hash function should not be used for any cryptographic purpose.
struct fnv1a_hasher
{
// Creates an empty hash object
fnv1a_hasher() = default;
~fnv1a_hasher() = default;
fnv1a_hasher(const fnv1a_hasher&) = default;
fnv1a_hasher(fnv1a_hasher&&) noexcept = default;
fnv1a_hasher& operator=(const fnv1a_hasher&) = default;
fnv1a_hasher& operator=(fnv1a_hasher&&) noexcept = default;
// Hashes a numeric value.
template <typename Tp, typename std::enable_if_t<std::is_arithmetic<Tp>::value, bool> = true>
void update(Tp data);
// Using the loop instead of "update(str, strlen(str))" to avoid looping twice
void update(const char* str);
// Hashes a byte array.
void update(const char* data, size_t size);
void update(std::string_view s) { update(s.data(), s.size()); }
void update(const std::string& s) { update(s.data(), s.size()); }
// Usage: uint64_t hashed_value = Hash::combine(33, false, "ABC", 458L, 3u, 'x');
template <typename Tp, typename... TailT>
static uint64_t combine(Tp&& arg, TailT&&... args);
// fnv1a_hasher.update_all(33, false, "ABC")` is shorthand for calling fnv1a_hasher.update(...)
// for each value in the same order
template <typename Tp, typename... TailT>
void update_all(Tp&& arg, TailT&&... args);
uint64_t digest() const { return m_result; }
private:
static constexpr uint64_t kFnv1a64OffsetBasis = 0xcbf29ce484222325;
static constexpr uint64_t kFnv1a64Prime = 0x100000001b3;
uint64_t m_result = kFnv1a64OffsetBasis;
};
template <typename Tp, typename std::enable_if_t<std::is_arithmetic<Tp>::value, bool>>
void
fnv1a_hasher::update(Tp data)
{
update(reinterpret_cast<const char*>(&data), sizeof(data));
}
inline void
fnv1a_hasher::update(const char* str)
{
constexpr auto max_n = 4096;
size_t n = 0;
for(const auto* p = str; *p != 0; ++p)
{
update(*p);
if(++n >= max_n) break; // prevent infinite loop
}
}
inline void
fnv1a_hasher::update(const char* data, size_t size)
{
for(size_t i = 0; i < size; ++i)
{
m_result ^= static_cast<uint8_t>(data[i]);
// Note: Arithmetic overflow of unsigned integers is well defined in C++ standard unlike
// signed integers.
m_result *= kFnv1a64Prime;
}
}
template <typename Tp, typename... TailT>
uint64_t
fnv1a_hasher::combine(Tp&& arg, TailT&&... args)
{
auto _hasher = fnv1a_hasher{};
_hasher.update_all(std::forward<Tp>(arg), std::forward<TailT>(args)...);
return _hasher.digest();
}
template <typename Tp, typename... TailT>
void
fnv1a_hasher::update_all(Tp&& arg, TailT&&... args)
{
update(arg);
if constexpr(sizeof...(TailT) > 0) update_all(std::forward<TailT>(args)...);
}
} // namespace common
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/common/md5sum.cpp
+378
Просмотреть файл
@@ -0,0 +1,378 @@
// MIT License
//
// Copyright (c) 2025 Advanced Micro Devices, Inc.
//
// 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/common/md5sum.hpp"
#include <cstdio>
#include <iomanip>
#include <string_view>
namespace rocprofiler
{
namespace common
{
namespace
{
using size_type = typename md5sum::size_type;
// Constants for md5sumTransform routine.
constexpr uint32_t S11 = 7;
constexpr uint32_t S12 = 12;
constexpr uint32_t S13 = 17;
constexpr uint32_t S14 = 22;
constexpr uint32_t S21 = 5;
constexpr uint32_t S22 = 9;
constexpr uint32_t S23 = 14;
constexpr uint32_t S24 = 20;
constexpr uint32_t S31 = 4;
constexpr uint32_t S32 = 11;
constexpr uint32_t S33 = 16;
constexpr uint32_t S34 = 23;
constexpr uint32_t S41 = 6;
constexpr uint32_t S42 = 10;
constexpr uint32_t S43 = 15;
constexpr uint32_t S44 = 21;
// low level logic operations
static inline uint32_t
F(uint32_t x, uint32_t y, uint32_t z);
static inline uint32_t
G(uint32_t x, uint32_t y, uint32_t z);
static inline uint32_t
H(uint32_t x, uint32_t y, uint32_t z);
static inline uint32_t
I(uint32_t x, uint32_t y, uint32_t z);
static inline uint32_t
rotate_left(uint32_t x, int n);
static inline void
FF(uint32_t& a, uint32_t b, uint32_t c, uint32_t d, uint32_t x, uint32_t s, uint32_t ac);
static inline void
GG(uint32_t& a, uint32_t b, uint32_t c, uint32_t d, uint32_t x, uint32_t s, uint32_t ac);
static inline void
HH(uint32_t& a, uint32_t b, uint32_t c, uint32_t d, uint32_t x, uint32_t s, uint32_t ac);
static inline void
II(uint32_t& a, uint32_t b, uint32_t c, uint32_t d, uint32_t x, uint32_t s, uint32_t ac);
// F, G, H and I are basic md5sum functions.
inline uint32_t
F(uint32_t x, uint32_t y, uint32_t z)
{
return (x & y) | (~x & z);
}
inline uint32_t
G(uint32_t x, uint32_t y, uint32_t z)
{
return (x & z) | (y & ~z);
}
inline uint32_t
H(uint32_t x, uint32_t y, uint32_t z)
{
return x ^ y ^ z;
}
inline uint32_t
I(uint32_t x, uint32_t y, uint32_t z)
{
return y ^ (x | ~z);
}
// rotate_left rotates x left n bits.
inline uint32_t
rotate_left(uint32_t x, int n)
{
return (x << n) | (x >> (32 - n));
}
// FF, GG, HH, and II transformations for rounds 1, 2, 3, and 4.
// Rotation is separate from addition to prevent recomputation.
inline void
FF(uint32_t& a, uint32_t b, uint32_t c, uint32_t d, uint32_t x, uint32_t s, uint32_t ac)
{
a = rotate_left(a + F(b, c, d) + x + ac, s) + b;
}
inline void
GG(uint32_t& a, uint32_t b, uint32_t c, uint32_t d, uint32_t x, uint32_t s, uint32_t ac)
{
a = rotate_left(a + G(b, c, d) + x + ac, s) + b;
}
inline void
HH(uint32_t& a, uint32_t b, uint32_t c, uint32_t d, uint32_t x, uint32_t s, uint32_t ac)
{
a = rotate_left(a + H(b, c, d) + x + ac, s) + b;
}
inline void
II(uint32_t& a, uint32_t b, uint32_t c, uint32_t d, uint32_t x, uint32_t s, uint32_t ac)
{
a = rotate_left(a + I(b, c, d) + x + ac, s) + b;
}
// decodes input (unsigned char) into output (uint32_t). Assumes len is a multiple of 4.
void
decode(uint32_t output[], const uint8_t input[], size_type len)
{
for(unsigned int i = 0, j = 0; j < len; i++, j += 4)
output[i] = ((uint32_t) input[j]) | (((uint32_t) input[j + 1]) << 8) |
(((uint32_t) input[j + 2]) << 16) | (((uint32_t) input[j + 3]) << 24);
}
// encodes input (uint32_t) into output (unsigned char). Assumes len is
// a multiple of 4.
void
encode(uint8_t output[], const uint32_t input[], size_type len)
{
for(size_type i = 0, j = 0; j < len; i++, j += 4)
{
output[j] = input[i] & 0xff;
output[j + 1] = (input[i] >> 8) & 0xff;
output[j + 2] = (input[i] >> 16) & 0xff;
output[j + 3] = (input[i] >> 24) & 0xff;
}
}
} // namespace
// apply md5sum algo on a block
void
md5sum::transform(const uint8_t block[blocksize])
{
uint32_t a = state[0], b = state[1], c = state[2], d = state[3], x[16];
decode(x, block, blocksize);
/* Round 1 */
FF(a, b, c, d, x[0], S11, 0xd76aa478); /* 1 */
FF(d, a, b, c, x[1], S12, 0xe8c7b756); /* 2 */
FF(c, d, a, b, x[2], S13, 0x242070db); /* 3 */
FF(b, c, d, a, x[3], S14, 0xc1bdceee); /* 4 */
FF(a, b, c, d, x[4], S11, 0xf57c0faf); /* 5 */
FF(d, a, b, c, x[5], S12, 0x4787c62a); /* 6 */
FF(c, d, a, b, x[6], S13, 0xa8304613); /* 7 */
FF(b, c, d, a, x[7], S14, 0xfd469501); /* 8 */
FF(a, b, c, d, x[8], S11, 0x698098d8); /* 9 */
FF(d, a, b, c, x[9], S12, 0x8b44f7af); /* 10 */
FF(c, d, a, b, x[10], S13, 0xffff5bb1); /* 11 */
FF(b, c, d, a, x[11], S14, 0x895cd7be); /* 12 */
FF(a, b, c, d, x[12], S11, 0x6b901122); /* 13 */
FF(d, a, b, c, x[13], S12, 0xfd987193); /* 14 */
FF(c, d, a, b, x[14], S13, 0xa679438e); /* 15 */
FF(b, c, d, a, x[15], S14, 0x49b40821); /* 16 */
/* Round 2 */
GG(a, b, c, d, x[1], S21, 0xf61e2562); /* 17 */
GG(d, a, b, c, x[6], S22, 0xc040b340); /* 18 */
GG(c, d, a, b, x[11], S23, 0x265e5a51); /* 19 */
GG(b, c, d, a, x[0], S24, 0xe9b6c7aa); /* 20 */
GG(a, b, c, d, x[5], S21, 0xd62f105d); /* 21 */
GG(d, a, b, c, x[10], S22, 0x2441453); /* 22 */
GG(c, d, a, b, x[15], S23, 0xd8a1e681); /* 23 */
GG(b, c, d, a, x[4], S24, 0xe7d3fbc8); /* 24 */
GG(a, b, c, d, x[9], S21, 0x21e1cde6); /* 25 */
GG(d, a, b, c, x[14], S22, 0xc33707d6); /* 26 */
GG(c, d, a, b, x[3], S23, 0xf4d50d87); /* 27 */
GG(b, c, d, a, x[8], S24, 0x455a14ed); /* 28 */
GG(a, b, c, d, x[13], S21, 0xa9e3e905); /* 29 */
GG(d, a, b, c, x[2], S22, 0xfcefa3f8); /* 30 */
GG(c, d, a, b, x[7], S23, 0x676f02d9); /* 31 */
GG(b, c, d, a, x[12], S24, 0x8d2a4c8a); /* 32 */
/* Round 3 */
HH(a, b, c, d, x[5], S31, 0xfffa3942); /* 33 */
HH(d, a, b, c, x[8], S32, 0x8771f681); /* 34 */
HH(c, d, a, b, x[11], S33, 0x6d9d6122); /* 35 */
HH(b, c, d, a, x[14], S34, 0xfde5380c); /* 36 */
HH(a, b, c, d, x[1], S31, 0xa4beea44); /* 37 */
HH(d, a, b, c, x[4], S32, 0x4bdecfa9); /* 38 */
HH(c, d, a, b, x[7], S33, 0xf6bb4b60); /* 39 */
HH(b, c, d, a, x[10], S34, 0xbebfbc70); /* 40 */
HH(a, b, c, d, x[13], S31, 0x289b7ec6); /* 41 */
HH(d, a, b, c, x[0], S32, 0xeaa127fa); /* 42 */
HH(c, d, a, b, x[3], S33, 0xd4ef3085); /* 43 */
HH(b, c, d, a, x[6], S34, 0x4881d05); /* 44 */
HH(a, b, c, d, x[9], S31, 0xd9d4d039); /* 45 */
HH(d, a, b, c, x[12], S32, 0xe6db99e5); /* 46 */
HH(c, d, a, b, x[15], S33, 0x1fa27cf8); /* 47 */
HH(b, c, d, a, x[2], S34, 0xc4ac5665); /* 48 */
/* Round 4 */
II(a, b, c, d, x[0], S41, 0xf4292244); /* 49 */
II(d, a, b, c, x[7], S42, 0x432aff97); /* 50 */
II(c, d, a, b, x[14], S43, 0xab9423a7); /* 51 */
II(b, c, d, a, x[5], S44, 0xfc93a039); /* 52 */
II(a, b, c, d, x[12], S41, 0x655b59c3); /* 53 */
II(d, a, b, c, x[3], S42, 0x8f0ccc92); /* 54 */
II(c, d, a, b, x[10], S43, 0xffeff47d); /* 55 */
II(b, c, d, a, x[1], S44, 0x85845dd1); /* 56 */
II(a, b, c, d, x[8], S41, 0x6fa87e4f); /* 57 */
II(d, a, b, c, x[15], S42, 0xfe2ce6e0); /* 58 */
II(c, d, a, b, x[6], S43, 0xa3014314); /* 59 */
II(b, c, d, a, x[13], S44, 0x4e0811a1); /* 60 */
II(a, b, c, d, x[4], S41, 0xf7537e82); /* 61 */
II(d, a, b, c, x[11], S42, 0xbd3af235); /* 62 */
II(c, d, a, b, x[2], S43, 0x2ad7d2bb); /* 63 */
II(b, c, d, a, x[9], S44, 0xeb86d391); /* 64 */
state[0] += a;
state[1] += b;
state[2] += c;
state[3] += d;
// Zeroize sensitive information.
memset(x, 0, sizeof x);
}
md5sum&
md5sum::update(std::string_view inp)
{
return update(inp.data(), inp.length());
}
// md5sum block update operation. Continues an md5sum message-digest
// operation, processing another message block
md5sum&
md5sum::update(const unsigned char input[], size_type length)
{
// compute number of bytes mod 64
size_type index = count[0] / 8 % blocksize;
// Update number of bits
if((count[0] += (length << 3)) < (length << 3)) count[1]++;
count[1] += (length >> 29);
// number of bytes we need to fill in buffer
size_type firstpart = 64 - index;
size_type i = 0;
// transform as many times as possible.
if(length >= firstpart)
{
// fill buffer first, transform
memcpy(&buffer[index], input, firstpart);
transform(buffer.data());
// transform chunks of blocksize (64 bytes)
for(i = firstpart; i + blocksize <= length; i += blocksize)
transform(&input[i]);
index = 0;
}
// buffer remaining input
memcpy(&buffer[index], &input[i], length - i);
return *this;
}
// for convenience provide a verson with signed char
md5sum&
md5sum::update(const char input[], size_type length)
{
return update((const unsigned char*) input, length);
}
// md5sum finalization. Ends an md5sum message-digest operation, writing the
// the message digest and zeroizing the context.
md5sum&
md5sum::finalize()
{
static unsigned char padding[64] = {0x80, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
if(!finalized)
{
// Save number of bits
unsigned char bits[8];
encode(bits, count.data(), 8);
// pad out to 56 mod 64.
size_type index = count[0] / 8 % 64;
size_type padLen = (index < 56) ? (56 - index) : (120 - index);
update(padding, padLen);
// Append length (before padding)
update(bits, 8);
// Store state in digest
encode(digest.data(), state.data(), 16);
// Zeroize sensitive information.
memset(buffer.data(), 0, sizeof buffer);
memset(count.data(), 0, sizeof count);
finalized = true;
}
return *this;
}
// return hex representation of digest as string
std::string
md5sum::hexdigest() const
{
if(!finalized) return std::string{};
char buf[33];
for(int i = 0; i < 16; i++)
snprintf(buf + i * 2, 3, "%02x", digest[i]);
buf[32] = '\0';
return std::string(buf);
}
std::string
md5sum::hexliteral() const
{
if(!finalized) return std::string{};
auto _oss = std::ostringstream{};
_oss << "X'";
for(auto itr : rawdigest())
_oss << std::hex << std::setw(2) << std::setfill('0') << static_cast<int>(itr);
_oss << "'";
return _oss.str();
}
std::ostream&
operator<<(std::ostream& out, md5sum md5)
{
return out << md5.hexdigest();
}
std::string
compute_md5sum(std::string_view inp)
{
return md5sum{inp}.finalize().hexdigest();
}
} // namespace common
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/common/md5sum.hpp
+138
Просмотреть файл
@@ -0,0 +1,138 @@
// MIT License
//
// Copyright (c) 2025 Advanced Micro Devices, Inc.
//
// 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 "lib/common/defines.hpp"
#include "lib/common/mpl.hpp"
#include <array>
#include <cstdint>
#include <cstring>
#include <iostream>
#include <sstream>
#include <string>
#include <type_traits>
namespace rocprofiler
{
namespace common
{
// helper function
std::string
compute_md5sum(std::string_view inp);
// helper function for array of string, string_view, or const char*
template <template <typename, typename...> class ContainerT, typename Tp, typename... TailT>
std::string
compute_md5sum(const ContainerT<Tp, TailT...>& inp,
std::enable_if_t<mpl::is_string_type<Tp>::value, int> = 0);
// a small class for calculating MD5 hashes of strings or byte arrays
//
// usage:
// 1) feed it blocks of uchars with update(...)
// 2) finalize()
// 3) get hexdigest() string
// or
// md5sum{...}.hexdigest()
//
// assumes that char is 8 bit and int is 32 bit
class md5sum
{
public:
using size_type = uint32_t; // must be 32bit
using raw_digest_t = std::array<uint8_t, 16>;
static constexpr int blocksize = 64;
template <typename Tp, typename... Args>
explicit md5sum(Tp&& arg, Args&&... args);
md5sum() = default;
~md5sum() = default;
md5sum(const md5sum&) = default;
md5sum(md5sum&&) = default;
md5sum& operator=(const md5sum&) = default;
md5sum& operator=(md5sum&&) = default;
md5sum& update(std::string_view inp);
md5sum& update(const unsigned char* buf, size_type length);
md5sum& update(const char* buf, size_type length);
md5sum& finalize();
std::string hexdigest() const;
std::string hexliteral() const;
raw_digest_t rawdigest() const { return digest; }
template <typename Tp, typename Up = std::enable_if_t<std::is_arithmetic<Tp>::value, int>>
md5sum& update(Tp inp);
friend std::ostream& operator<<(std::ostream&, md5sum md5);
private:
void transform(const uint8_t block[blocksize]);
bool finalized = false;
// 64bit counter for number of bits (lo, hi)
std::array<uint32_t, 2> count = {0, 0};
std::array<uint8_t, blocksize> buffer{}; // overflow bytes from last 64 byte chunk
// digest so far, initialized to magic initialization constants.
std::array<uint32_t, 4> state = {0x67452301, 0xefcdab89, 0x98badcfe, 0x10325476};
std::array<uint8_t, 16> digest{}; // result
};
template <typename Tp, typename... Args>
md5sum::md5sum(Tp&& arg, Args&&... args)
{
auto _update = [&](auto&& _val) {
using value_type = common::mpl::unqualified_type_t<decltype(_val)>;
static_assert(!std::is_pointer<value_type>::value,
"constructor cannot be called with pointer argument");
update(std::forward<decltype(_val)>(_val));
};
_update(std::forward<Tp>(arg));
(_update(std::forward<Args>(args)), ...);
finalize();
}
template <typename Tp, typename Up>
md5sum&
md5sum::update(Tp inp)
{
static_assert(std::is_arithmetic<Tp>::value, "expected arithmetic type");
return update(reinterpret_cast<const char*>(&inp), sizeof(Tp));
}
template <template <typename, typename...> class ContainerT, typename Tp, typename... TailT>
std::string
compute_md5sum(const ContainerT<Tp, TailT...>& inp,
std::enable_if_t<mpl::is_string_type<Tp>::value, int>)
{
auto _val = md5sum{};
for(const auto& itr : inp)
_val.update(std::string_view{inp});
_val.finalize();
return _val.hexdigest();
}
} // namespace common
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/common/mpl.hpp
+1 -4
Просмотреть файл
@@ -38,10 +38,7 @@ namespace mpl
{
// dummy tuple with low instantiation cost
template <typename... Tp>
struct type_list
{
static constexpr auto size() { return sizeof...(Tp); }
};
using type_list = impl::type_list<Tp...>;
/// get the index of a type in expansion
template <typename Tp, typename Type>
projects/rocprofiler-sdk/source/lib/common/simple_timer.cpp
+93
Просмотреть файл
@@ -0,0 +1,93 @@
// MIT License
//
// Copyright (c) 2025 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/common/simple_timer.hpp"
#include "lib/common/logging.hpp"
#include <fmt/format.h>
#include <chrono>
#include <cstdint>
#include <string>
#include <string_view>
#include <type_traits>
namespace rocprofiler
{
namespace common
{
simple_timer::simple_timer(std::string&& label)
: m_label{std::move(label)}
{
start();
}
simple_timer::simple_timer(std::string&& label, defer_start)
: m_label{std::move(label)}
{}
simple_timer::~simple_timer()
{
if(m_quiet)
return;
else if(m_end <= m_beg)
stop();
ROCP_WARNING << fmt::format("{} :: {:12.6f} sec", m_label, get());
}
void
simple_timer::start()
{
m_beg = clock_type::now();
}
void
simple_timer::stop()
{
m_end = clock_type::now();
}
double
simple_timer::get() const
{
if(m_end <= m_beg) return {};
return std::chrono::duration_cast<std::chrono::duration<double>>(m_end - m_beg).count();
}
size_t
simple_timer::get_nsec() const
{
if(m_end <= m_beg) return {};
return std::chrono::duration_cast<std::chrono::nanoseconds>(m_end - m_beg).count();
}
std::ostream&
operator<<(std::ostream& _os, const simple_timer& _val)
{
_val.set_quiet(true);
_os << fmt::format("{} :: {:12.6f} sec", _val.label(), _val.get());
return _os;
}
} // namespace common
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/common/simple_timer.hpp
+67
Просмотреть файл
@@ -0,0 +1,67 @@
// MIT License
//
// Copyright (c) 2025 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 "lib/common/defines.hpp"
#include <chrono>
#include <cstdint>
#include <string>
#include <string_view>
#include <type_traits>
namespace rocprofiler
{
namespace common
{
struct defer_start
{};
struct simple_timer
{
using duration_t = std::chrono::duration<double>;
explicit simple_timer(std::string&& label);
explicit simple_timer(std::string&& label, defer_start);
~simple_timer();
void start();
void stop();
double get() const;
size_t get_nsec() const;
std::string_view label() const { return std::string_view{m_label}; }
void set_quiet(bool v) const { m_quiet = v; }
friend std::ostream& operator<<(std::ostream& _os, const simple_timer& _val);
private:
using clock_type = std::chrono::steady_clock;
using time_point_t = std::chrono::time_point<clock_type, std::chrono::nanoseconds>;
std::string m_label = {};
time_point_t m_beg = {};
time_point_t m_end = {};
mutable bool m_quiet = false;
};
} // namespace common
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/common/static_tl_object.cpp
+78
Просмотреть файл
@@ -0,0 +1,78 @@
// MIT License
//
// Copyright (c) 2023 Advanced Micro Devices, Inc.
//
// 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/common/static_tl_object.hpp"
#include "lib/common/scope_destructor.hpp"
#include <array>
#include <cstddef>
#include <functional>
#include <mutex>
#include <stack>
namespace rocprofiler
{
namespace common
{
namespace
{
auto*&
get_static_tl_object_stack()
{
static thread_local auto* _v = new std::stack<static_dtor_func_t>{};
return _v;
}
} // namespace
thread_local auto _static_tl_object_dtor = std::optional<scope_destructor>{};
void
destroy_static_tl_objects()
{
if(auto*& _stack = get_static_tl_object_stack(); _stack)
{
while(!_stack->empty())
{
auto& itr = _stack->top();
if(itr) itr();
_stack->pop();
}
delete _stack;
_stack = nullptr;
}
}
void
register_static_tl_dtor(static_dtor_func_t&& _func)
{
// make sure the thread-local scope destructor exists
if(!_static_tl_object_dtor)
_static_tl_object_dtor = scope_destructor{[]() { destroy_static_tl_objects(); }};
if(auto*& _stack = get_static_tl_object_stack(); _stack)
{
_stack->push(_func);
}
}
} // namespace common
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/common/static_tl_object.hpp
+4 -1
Просмотреть файл
@@ -39,6 +39,9 @@ namespace common
void
destroy_static_tl_objects();
void
register_static_tl_dtor(static_dtor_func_t&&);
/**
* @brief This struct is used to create static singleton objects which have the properties of a
* heap-allocated static object without a memory leak.
@@ -97,7 +100,7 @@ static_tl_object<Tp, ContextT>::construct(Args&&... args)
{
static thread_local auto _once = std::once_flag{};
std::call_once(_once, []() {
register_static_dtor([]() {
register_static_tl_dtor([]() {
if(static_tl_object<Tp, ContextT>::m_object)
{
static_tl_object<Tp, ContextT>::m_object->~Tp();
projects/rocprofiler-sdk/source/lib/common/utility.cpp
+10 -1
Просмотреть файл
@@ -22,6 +22,7 @@
//
#include "lib/common/utility.hpp"
#include "lib/common/defines.hpp"
#include "lib/common/logging.hpp"
#include <unistd.h>
@@ -32,7 +33,6 @@
#include <sstream>
#include <string>
#include <vector>
#include "lib/common/defines.hpp"
namespace rocprofiler
{
@@ -63,6 +63,8 @@ get_clock_name(clockid_t _id)
return "CLOCK_UNKNOWN";
#undef CLOCK_NAME_CASE_STATEMENT
}
auto _process_init_ns = timestamp_ns();
} // namespace
uint64_t
@@ -89,6 +91,13 @@ get_clock_period_ns_impl(clockid_t _clk_id)
return (static_cast<uint64_t>(ts.tv_sec) * nanosec) + static_cast<uint64_t>(ts.tv_nsec);
}
uint64_t
get_process_start_time_ns(pid_t _pid)
{
if(_pid == getpid()) return _process_init_ns;
return 0;
}
std::vector<std::string>
read_command_line(pid_t _pid)
{
projects/rocprofiler-sdk/source/lib/common/utility.hpp
+4
Просмотреть файл
@@ -94,6 +94,10 @@ timestamp_ns()
return get_ticks(_clk) / _clk_period;
}
// returns the process start time (in CLOCK_BOOTTIME nanoseconds) via /proc/<pid>/stat
uint64_t
get_process_start_time_ns(pid_t _pid);
std::vector<std::string>
read_command_line(pid_t _pid);
projects/rocprofiler-sdk/source/lib/output/CMakeLists.txt
+4 -1
Просмотреть файл
@@ -21,6 +21,7 @@ set(TOOL_OUTPUT_HEADERS
kernel_symbol_info.hpp
host_symbol_info.hpp
metadata.hpp
node_info.hpp
output_config.hpp
output_key.hpp
output_stream.hpp
@@ -41,6 +42,7 @@ set(TOOL_OUTPUT_SOURCES
generatePerfetto.cpp
generateStats.cpp
metadata.cpp
node_info.cpp
output_config.cpp
output_key.cpp
output_stream.cpp
@@ -64,4 +66,5 @@ target_link_libraries(
rocprofiler-sdk::rocprofiler-sdk-otf2
rocprofiler-sdk::rocprofiler-sdk-amd-comgr
rocprofiler-sdk::rocprofiler-sdk-dw
rocprofiler-sdk::rocprofiler-sdk-elf)
rocprofiler-sdk::rocprofiler-sdk-elf
rocprofiler-sdk::rocprofiler-sdk-sqlite3)
projects/rocprofiler-sdk/source/lib/output/agent_info.hpp
+9 -3
Просмотреть файл
@@ -41,6 +41,7 @@ struct agent_info : rocprofiler_agent_v0_t
: base_type{_base}
{}
agent_info() = default;
~agent_info() = default;
agent_info(const agent_info&) = default;
agent_info(agent_info&&) noexcept = default;
@@ -58,14 +59,19 @@ using agent_info_map_t = std::unordered_map<rocprofiler_agent_id_t, agent_info>;
namespace cereal
{
#define SAVE_DATA_FIELD(FIELD) ar(make_nvp(#FIELD, data.FIELD))
template <typename ArchiveT>
void
save(ArchiveT& ar, const ::rocprofiler::tool::agent_info& data)
{
cereal::save(ar, static_cast<const rocprofiler_agent_v0_t&>(data));
ar(cereal::make_nvp("gpu_index", data.gpu_index));
}
#undef SAVE_DATA_FIELD
template <typename ArchiveT>
void
load(ArchiveT& ar, ::rocprofiler::tool::agent_info& data)
{
cereal::load(ar, static_cast<rocprofiler_agent_v0_t&>(data));
ar(cereal::make_nvp("gpu_index", data.gpu_index));
}
} // namespace cereal
projects/rocprofiler-sdk/source/lib/output/buffered_output.hpp
+13 -6
Просмотреть файл
@@ -63,9 +63,9 @@ struct buffered_output
void clear();
void destroy();
uint64_t get_num_bytes() const;
generator<Tp> get_generator() const { return generator<Tp>{get_tmp_file_buffer<Tp>(DomainT)}; }
std::deque<Tp> load_all();
uint64_t get_num_bytes() const;
file_generator<Tp> get_generator() const;
std::deque<Tp> load_all();
stats_entry_t stats = {};
@@ -98,6 +98,13 @@ buffered_output<Tp, DomainT>::read()
read_tmp_file<type>(buffer_type_v);
}
template <typename Tp, domain_type DomainT>
file_generator<Tp>
buffered_output<Tp, DomainT>::get_generator() const
{
return file_generator<Tp>{get_tmp_file_buffer<Tp>(DomainT)};
}
template <typename Tp, domain_type DomainT>
std::deque<Tp>
buffered_output<Tp, DomainT>::load_all()
@@ -169,7 +176,7 @@ using scratch_memory_buffered_output_t =
buffered_output<rocprofiler_buffer_tracing_scratch_memory_record_t,
domain_type::SCRATCH_MEMORY>;
using memory_allocation_buffered_output_t =
buffered_output<rocprofiler_buffer_tracing_memory_allocation_record_t,
buffered_output<tool_buffer_tracing_memory_allocation_ext_record_t,
domain_type::MEMORY_ALLOCATION>;
using counter_records_buffered_output_t =
::rocprofiler::tool::buffered_output<rocprofiler::tool::serialized_counter_record_t,
@@ -181,9 +188,9 @@ using rocdecode_buffered_output_t =
buffered_output<rocprofiler_buffer_tracing_rocdecode_api_ext_record_t, domain_type::ROCDECODE>;
using rocjpeg_buffered_output_t =
buffered_output<rocprofiler_buffer_tracing_rocjpeg_api_record_t, domain_type::ROCJPEG>;
using kernel_dispatch_buffered_output_with_stream_t =
using kernel_dispatch_buffered_output_ext_t =
buffered_output<tool_buffer_tracing_kernel_dispatch_ext_record_t, domain_type::KERNEL_DISPATCH>;
using memory_copy_buffered_output_with_stream_t =
using memory_copy_buffered_output_ext_t =
buffered_output<tool_buffer_tracing_memory_copy_ext_record_t, domain_type::MEMORY_COPY>;
using pc_sampling_stochastic_buffered_output_t =
buffered_output<rocprofiler::tool::rocprofiler_tool_pc_sampling_stochastic_record_t,
projects/rocprofiler-sdk/source/lib/output/counter_info.hpp
+6 -7
Просмотреть файл
@@ -41,7 +41,7 @@ namespace tool
constexpr uint32_t lds_block_size = 128 * 4;
using counter_dimension_id_vec_t = std::vector<rocprofiler_counter_dimension_id_t>;
using counter_dimension_info_vec_t = std::vector<rocprofiler_record_dimension_info_t>;
using counter_dimension_info_vec_t = std::vector<rocprofiler_counter_record_dimension_info_t>;
struct tool_counter_info : rocprofiler_counter_info_v1_t
{
@@ -93,10 +93,10 @@ struct tool_counter_record_t
{
using container_type = std::vector<tool_counter_value_t>;
uint64_t thread_id = 0;
rocprofiler_dispatch_counting_service_data_t dispatch_data = {};
serialized_counter_record_t record = {};
uint64_t kernel_rename_val = {};
uint64_t thread_id = 0;
rocprofiler_dispatch_counting_service_data_t dispatch_data = {};
serialized_counter_record_t record = {};
rocprofiler_stream_id_t stream_id = {.handle = 0};
template <typename ArchiveT>
void save(ArchiveT& ar) const
@@ -107,7 +107,7 @@ struct tool_counter_record_t
ar(cereal::make_nvp("thread_id", thread_id));
ar(cereal::make_nvp("dispatch_data", dispatch_data));
ar(cereal::make_nvp("records", tmp));
ar(cereal::make_nvp("kernel_rename_val", kernel_rename_val));
ar(cereal::make_nvp("stream_id", stream_id));
}
container_type read() const;
@@ -126,7 +126,6 @@ save(ArchiveT& ar, const ::rocprofiler::tool::tool_counter_info& data)
{
SAVE_DATA_FIELD(agent_id);
cereal::save(ar, static_cast<const rocprofiler_counter_info_v1_t&>(data));
SAVE_DATA_FIELD(dimension_ids);
}
#undef SAVE_DATA_FIELD
projects/rocprofiler-sdk/source/lib/output/format_path.cpp
+29 -10
Просмотреть файл
@@ -82,7 +82,8 @@ const auto env_regexes =
std::string
format_path_impl(std::string _fpath, const std::vector<output_key>& _keys)
{
if(_fpath.find('%') == std::string::npos && _fpath.find('$') == std::string::npos)
if(_fpath.find('%') == std::string::npos && _fpath.find('{') == std::string::npos &&
_fpath.find('$') == std::string::npos)
return _fpath;
auto _replace = [](auto& _v, const output_key& pitr) {
@@ -133,9 +134,9 @@ format_path_impl(std::string _fpath, const std::vector<output_key>& _keys)
// remove %arg<N>% where N >= argc
try
{
std::regex _re{"(.*)%(arg[0-9]+)%([-/_]*)(.*)"};
auto _re = std::regex{"(.*)(%|\\{)(arg[0-9]+)(%|\\})([-/_]*)(.*)"};
while(std::regex_search(_fpath, _re))
_fpath = std::regex_replace(_fpath, _re, "$1$4");
_fpath = std::regex_replace(_fpath, _re, "$1$6");
} catch(std::exception& _e)
{
ROCP_WARNING << "[rocprofiler] " << __FUNCTION__ << " threw an exception :: " << _e.what()
@@ -148,7 +149,8 @@ format_path_impl(std::string _fpath, const std::vector<output_key>& _keys)
std::string
format_path(std::string&& _fpath, const std::vector<output_key>& _keys)
{
if(_fpath.find('%') == std::string::npos && _fpath.find('$') == std::string::npos)
if(_fpath.find('%') == std::string::npos && _fpath.find('{') == std::string::npos &&
_fpath.find('$') == std::string::npos)
return _fpath;
auto _ref = _fpath;
@@ -156,23 +158,40 @@ format_path(std::string&& _fpath, const std::vector<output_key>& _keys)
return (_fpath == _ref) ? _fpath : format_path(std::move(_fpath), _keys);
}
template <typename Tp>
Tp
get_variable_env(Tp _default_v, std::initializer_list<std::string_view>&& _options)
{
// set env variables towards end override preceding environment variables
auto _val = _default_v;
for(auto itr : _options)
_val = common::get_env<Tp>(itr, std::move(_val));
return _val;
}
} // namespace
int
get_mpi_size()
{
static int _v = common::get_env<int>(
"OMPI_COMM_WORLD_SIZE",
common::get_env<int>("MV2_COMM_WORLD_SIZE", common::get_env<int>("MPI_SIZE", 0)));
static int _v = get_variable_env<int>(0,
{"MPI_SIZE", // most generic to most runtime-specific
"MPI_LOCALNRANKS",
"MPI_NRANKS",
"MV2_COMM_WORLD_SIZE",
"OMPI_COMM_WORLD_SIZE"});
return _v;
}
int
get_mpi_rank()
{
static int _v = common::get_env<int>(
"OMPI_COMM_WORLD_RANK",
common::get_env<int>("MV2_COMM_WORLD_RANK", common::get_env<int>("MPI_RANK", -1)));
static int _v = get_variable_env<int>(0,
{"MPI_RANK", // most generic to most runtime-specific
"MPI_LOCALRANKID",
"MPI_RANKID",
"MV2_COMM_WORLD_RANK",
"OMPI_COMM_WORLD_RANK"});
return _v;
}
projects/rocprofiler-sdk/source/lib/output/generateCSV.cpp
+6 -6
Просмотреть файл
@@ -287,7 +287,7 @@ generate_csv(const output_config&
{
auto row_ss = std::stringstream{};
auto kernel_name = tool_metadata.get_kernel_name(record.dispatch_info.kernel_id,
record.kernel_rename_val);
record.correlation_id.external.value);
rocprofiler::tool::csv::kernel_trace_with_stream_csv_encoder::write_row(
row_ss,
tool_metadata.get_kind_name(record.kind),
@@ -445,10 +445,10 @@ generate_csv(const output_config& cfg,
}
void
generate_csv(const output_config& cfg,
const metadata& tool_metadata,
const generator<rocprofiler_buffer_tracing_memory_allocation_record_t>& data,
const stats_entry_t& stats)
generate_csv(const output_config& cfg,
const metadata& tool_metadata,
const generator<tool_buffer_tracing_memory_allocation_ext_record_t>& data,
const stats_entry_t& stats)
{
if(data.empty()) return;
@@ -626,7 +626,7 @@ generate_csv(const output_config& cfg,
record.thread_id,
magnitude(record.dispatch_data.dispatch_info.grid_size),
record.dispatch_data.dispatch_info.kernel_id,
tool_metadata.get_kernel_name(kernel_id, record.kernel_rename_val),
tool_metadata.get_kernel_name(kernel_id, correlation_id.external.value),
magnitude(record.dispatch_data.dispatch_info.workgroup_size),
lds_block_size_v,
record.dispatch_data.dispatch_info.private_segment_size,
projects/rocprofiler-sdk/source/lib/output/generateCSV.hpp
+5 -4
Просмотреть файл
@@ -100,10 +100,11 @@ generate_csv(const output_config& c
const stats_entry_t& stats);
void
generate_csv(const output_config& cfg,
const metadata& tool_metadata,
const generator<rocprofiler_buffer_tracing_memory_allocation_record_t>& data,
const stats_entry_t& stats);
generate_csv(const output_config& cfg,
const metadata& tool_metadata,
const generator<tool_buffer_tracing_memory_allocation_ext_record_t>& data,
const stats_entry_t& stats);
void
generate_csv(const output_config& cfg,
const metadata& tool_metadata,
projects/rocprofiler-sdk/source/lib/output/generateJSON.cpp
+22 -23
Просмотреть файл
@@ -100,16 +100,14 @@ write_json(json_output& json_ar,
{
// metadata
{
auto timestamps =
timestamps_t{tool_metadata.process_start_ns, tool_metadata.process_end_ns};
json_ar.setNextName("metadata");
json_ar.startNode();
json_ar(cereal::make_nvp("node", tool_metadata.node_data));
json_ar(cereal::make_nvp("pid", pid));
json_ar(cereal::make_nvp("init_time", timestamps.app_start_time));
json_ar(cereal::make_nvp("fini_time", timestamps.app_end_time));
json_ar(cereal::make_nvp("init_time", tool_metadata.process_start_ns));
json_ar(cereal::make_nvp("fini_time", tool_metadata.process_end_ns));
json_ar(cereal::make_nvp("command", tool_metadata.command_line));
json_ar(cereal::make_nvp("config", cfg));
json_ar(cereal::make_nvp("command", common::read_command_line(pid)));
json_ar.finishNode();
}
@@ -183,23 +181,24 @@ write_json(json_output& json_ar,
}
void
write_json(json_output& json_ar,
const output_config& /*cfg*/,
const metadata& /*tool_metadata*/,
const domain_stats_vec_t& domain_stats,
generator<rocprofiler_buffer_tracing_hip_api_ext_record_t>&& hip_api_gen,
generator<rocprofiler_buffer_tracing_hsa_api_record_t> hsa_api_gen,
generator<tool_buffer_tracing_kernel_dispatch_ext_record_t> kernel_dispatch_gen,
generator<tool_buffer_tracing_memory_copy_ext_record_t> memory_copy_gen,
generator<tool_counter_record_t> counter_collection_gen,
generator<rocprofiler_buffer_tracing_marker_api_record_t> marker_api_gen,
generator<rocprofiler_buffer_tracing_scratch_memory_record_t> scratch_memory_gen,
generator<rocprofiler_buffer_tracing_rccl_api_record_t> rccl_api_gen,
generator<rocprofiler_buffer_tracing_memory_allocation_record_t> memory_allocation_gen,
generator<rocprofiler_buffer_tracing_rocdecode_api_ext_record_t> rocdecode_api_gen,
generator<rocprofiler_buffer_tracing_rocjpeg_api_record_t> rocjpeg_api_gen,
generator<rocprofiler_tool_pc_sampling_host_trap_record_t> pc_sampling_host_trap_gen,
generator<rocprofiler_tool_pc_sampling_stochastic_record_t> pc_sampling_stochastic_gen)
write_json(
json_output& json_ar,
const output_config& /*cfg*/,
const metadata& /*tool_metadata*/,
const domain_stats_vec_t& domain_stats,
const generator<rocprofiler_buffer_tracing_hip_api_ext_record_t>& hip_api_gen,
const generator<rocprofiler_buffer_tracing_hsa_api_record_t>& hsa_api_gen,
const generator<tool_buffer_tracing_kernel_dispatch_ext_record_t>& kernel_dispatch_gen,
const generator<tool_buffer_tracing_memory_copy_ext_record_t>& memory_copy_gen,
const generator<tool_counter_record_t>& counter_collection_gen,
const generator<rocprofiler_buffer_tracing_marker_api_record_t>& marker_api_gen,
const generator<rocprofiler_buffer_tracing_scratch_memory_record_t>& scratch_memory_gen,
const generator<rocprofiler_buffer_tracing_rccl_api_record_t>& rccl_api_gen,
const generator<tool_buffer_tracing_memory_allocation_ext_record_t>& memory_allocation_gen,
const generator<rocprofiler_buffer_tracing_rocdecode_api_ext_record_t>& rocdecode_api_gen,
const generator<rocprofiler_buffer_tracing_rocjpeg_api_record_t>& rocjpeg_api_gen,
const generator<rocprofiler_tool_pc_sampling_host_trap_record_t>& pc_sampling_host_trap_gen,
const generator<rocprofiler_tool_pc_sampling_stochastic_record_t>& pc_sampling_stochastic_gen)
{
// summary
{
projects/rocprofiler-sdk/source/lib/output/generateJSON.hpp
+18 -17
Просмотреть файл
@@ -81,22 +81,23 @@ void
write_json(json_output&, const output_config& cfg, const metadata& tool_metadata, uint64_t pid);
void
write_json(json_output& json_ar,
const output_config& cfg,
const metadata& tool_metadata,
const domain_stats_vec_t& domain_stats,
generator<rocprofiler_buffer_tracing_hip_api_ext_record_t>&& hip_api_gen,
generator<rocprofiler_buffer_tracing_hsa_api_record_t> hsa_api_gen,
generator<tool_buffer_tracing_kernel_dispatch_ext_record_t> kernel_dispatch_gen,
generator<tool_buffer_tracing_memory_copy_ext_record_t> memory_copy_gen,
generator<tool_counter_record_t> counter_collection_gen,
generator<rocprofiler_buffer_tracing_marker_api_record_t> marker_api_gen,
generator<rocprofiler_buffer_tracing_scratch_memory_record_t> scratch_memory_gen,
generator<rocprofiler_buffer_tracing_rccl_api_record_t> rccl_api_gen,
generator<rocprofiler_buffer_tracing_memory_allocation_record_t> memory_allocation_gen,
generator<rocprofiler_buffer_tracing_rocdecode_api_ext_record_t> rocdecode_api_gen,
generator<rocprofiler_buffer_tracing_rocjpeg_api_record_t> rocjpeg_api_gen,
generator<rocprofiler_tool_pc_sampling_host_trap_record_t> pc_sampling_host_trap_gen,
generator<rocprofiler_tool_pc_sampling_stochastic_record_t> pc_sampling_stochastic_gen);
write_json(
json_output& json_ar,
const output_config& cfg,
const metadata& tool_metadata,
const domain_stats_vec_t& domain_stats,
const generator<rocprofiler_buffer_tracing_hip_api_ext_record_t>& hip_api_gen,
const generator<rocprofiler_buffer_tracing_hsa_api_record_t>& hsa_api_gen,
const generator<tool_buffer_tracing_kernel_dispatch_ext_record_t>& kernel_dispatch_gen,
const generator<tool_buffer_tracing_memory_copy_ext_record_t>& memory_copy_gen,
const generator<tool_counter_record_t>& counter_collection_gen,
const generator<rocprofiler_buffer_tracing_marker_api_record_t>& marker_api_gen,
const generator<rocprofiler_buffer_tracing_scratch_memory_record_t>& scratch_memory_gen,
const generator<rocprofiler_buffer_tracing_rccl_api_record_t>& rccl_api_gen,
const generator<tool_buffer_tracing_memory_allocation_ext_record_t>& memory_allocation_gen,
const generator<rocprofiler_buffer_tracing_rocdecode_api_ext_record_t>& rocdecode_api_gen,
const generator<rocprofiler_buffer_tracing_rocjpeg_api_record_t>& rocjpeg_api_gen,
const generator<rocprofiler_tool_pc_sampling_host_trap_record_t>& pc_sampling_host_trap_gen,
const generator<rocprofiler_tool_pc_sampling_stochastic_record_t>& pc_sampling_stochastic_gen);
} // namespace tool
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/output/generateOTF2.cpp
+17 -17
Просмотреть файл
@@ -355,21 +355,20 @@ create_attribute_list()
} // namespace
void
write_otf2(
const output_config& cfg,
const metadata& tool_metadata,
uint64_t pid,
const std::vector<agent_info>& agent_data,
std::deque<rocprofiler_buffer_tracing_hip_api_ext_record_t>* hip_api_data,
std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>* hsa_api_data,
std::deque<tool_buffer_tracing_kernel_dispatch_ext_record_t>* kernel_dispatch_data,
std::deque<tool_buffer_tracing_memory_copy_ext_record_t>* memory_copy_data,
std::deque<rocprofiler_buffer_tracing_marker_api_record_t>* marker_api_data,
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* /*scratch_memory_data*/,
std::deque<rocprofiler_buffer_tracing_rccl_api_record_t>* rccl_api_data,
std::deque<rocprofiler_buffer_tracing_memory_allocation_record_t>* memory_allocation_data,
std::deque<rocprofiler_buffer_tracing_rocdecode_api_ext_record_t>* rocdecode_api_data,
std::deque<rocprofiler_buffer_tracing_rocjpeg_api_record_t>* rocjpeg_api_data)
write_otf2(const output_config& cfg,
const metadata& tool_metadata,
uint64_t pid,
const std::vector<agent_info>& agent_data,
std::deque<rocprofiler_buffer_tracing_hip_api_ext_record_t>* hip_api_data,
std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>* hsa_api_data,
std::deque<tool_buffer_tracing_kernel_dispatch_ext_record_t>* kernel_dispatch_data,
std::deque<tool_buffer_tracing_memory_copy_ext_record_t>* memory_copy_data,
std::deque<rocprofiler_buffer_tracing_marker_api_record_t>* marker_api_data,
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* /*scratch_memory_data*/,
std::deque<rocprofiler_buffer_tracing_rccl_api_record_t>* rccl_api_data,
std::deque<tool_buffer_tracing_memory_allocation_ext_record_t>* memory_allocation_data,
std::deque<rocprofiler_buffer_tracing_rocdecode_api_ext_record_t>* rocdecode_api_data,
std::deque<rocprofiler_buffer_tracing_rocjpeg_api_record_t>* rocjpeg_api_data)
{
namespace sdk = ::rocprofiler::sdk;
@@ -481,7 +480,7 @@ write_otf2(
tool_metadata.get_agent_index(_agent->id, cfg.agent_index_value);
evt.name = fmt::format("Thread {}, Copy to {} {}",
tid,
agent_index_info.type,
std::string{agent_index_info.type},
agent_index_info.as_string("-"));
}
}
@@ -697,7 +696,8 @@ write_otf2(
const auto* sym = _get_kernel_sym_data(info);
CHECK(sym != nullptr);
auto name = tool_metadata.get_kernel_name(info.kernel_id, itr.kernel_rename_val);
auto name =
tool_metadata.get_kernel_name(info.kernel_id, itr.correlation_id.external.value);
_hash_data.emplace(
get_hash_id(name),
region_info{std::string{name}, OTF2_REGION_ROLE_FUNCTION, OTF2_PARADIGM_HIP});
projects/rocprofiler-sdk/source/lib/output/generateOTF2.hpp
+14 -15
Просмотреть файл
@@ -35,20 +35,19 @@ namespace rocprofiler
namespace tool
{
void
write_otf2(
const output_config& cfg,
const metadata& tool_metadata,
uint64_t pid,
const std::vector<agent_info>& agent_data,
std::deque<rocprofiler_buffer_tracing_hip_api_ext_record_t>* hip_api_data,
std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>* hsa_api_data,
std::deque<tool_buffer_tracing_kernel_dispatch_ext_record_t>* kernel_dispatch_data,
std::deque<tool_buffer_tracing_memory_copy_ext_record_t>* memory_copy_data,
std::deque<rocprofiler_buffer_tracing_marker_api_record_t>* marker_api_data,
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* scratch_memory_data,
std::deque<rocprofiler_buffer_tracing_rccl_api_record_t>* rccl_api_data,
std::deque<rocprofiler_buffer_tracing_memory_allocation_record_t>* memory_allocation_data,
std::deque<rocprofiler_buffer_tracing_rocdecode_api_ext_record_t>* rocdecode_api_data,
std::deque<rocprofiler_buffer_tracing_rocjpeg_api_record_t>* rocjpeg_api_data);
write_otf2(const output_config& cfg,
const metadata& tool_metadata,
uint64_t pid,
const std::vector<agent_info>& agent_data,
std::deque<rocprofiler_buffer_tracing_hip_api_ext_record_t>* hip_api_data,
std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>* hsa_api_data,
std::deque<tool_buffer_tracing_kernel_dispatch_ext_record_t>* kernel_dispatch_data,
std::deque<tool_buffer_tracing_memory_copy_ext_record_t>* memory_copy_data,
std::deque<rocprofiler_buffer_tracing_marker_api_record_t>* marker_api_data,
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* scratch_memory_data,
std::deque<rocprofiler_buffer_tracing_rccl_api_record_t>* rccl_api_data,
std::deque<tool_buffer_tracing_memory_allocation_ext_record_t>* memory_allocation_data,
std::deque<rocprofiler_buffer_tracing_rocdecode_api_ext_record_t>* rocdecode_api_data,
std::deque<rocprofiler_buffer_tracing_rocjpeg_api_record_t>* rocjpeg_api_data);
} // namespace tool
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/output/generatePerfetto.cpp
+1 -1
Просмотреть файл
@@ -76,7 +76,7 @@ write_perfetto(
const generator<rocprofiler_buffer_tracing_marker_api_record_t>& marker_api_gen,
const generator<rocprofiler_buffer_tracing_scratch_memory_record_t>& /*scratch_memory_gen*/,
const generator<rocprofiler_buffer_tracing_rccl_api_record_t>& rccl_api_gen,
const generator<rocprofiler_buffer_tracing_memory_allocation_record_t>& memory_allocation_gen,
const generator<tool_buffer_tracing_memory_allocation_ext_record_t>& memory_allocation_gen,
const generator<rocprofiler_buffer_tracing_rocdecode_api_ext_record_t>& rocdecode_api_gen,
const generator<rocprofiler_buffer_tracing_rocjpeg_api_record_t>& rocjpeg_api_gen)
{
projects/rocprofiler-sdk/source/lib/output/generatePerfetto.hpp
+1 -1
Просмотреть файл
@@ -48,7 +48,7 @@ write_perfetto(
const generator<rocprofiler_buffer_tracing_marker_api_record_t>& marker_api_gen,
const generator<rocprofiler_buffer_tracing_scratch_memory_record_t>& scratch_memory_gen,
const generator<rocprofiler_buffer_tracing_rccl_api_record_t>& rccl_api_gen,
const generator<rocprofiler_buffer_tracing_memory_allocation_record_t>& memory_allocation_gen,
const generator<tool_buffer_tracing_memory_allocation_ext_record_t>& memory_allocation_gen,
const generator<rocprofiler_buffer_tracing_rocdecode_api_ext_record_t>& rocdecode_api_gen,
const generator<rocprofiler_buffer_tracing_rocjpeg_api_record_t>& rocjpeg_api_gen);
} // namespace tool
projects/rocprofiler-sdk/source/lib/output/generateStats.cpp
+3 -3
Просмотреть файл
@@ -72,7 +72,7 @@ generate_stats(const output_config& /*cfg*/,
for(auto record : data.get(ditr))
{
auto kernel_name = tool_metadata.get_kernel_name(record.dispatch_info.kernel_id,
record.kernel_rename_val);
record.correlation_id.external.value);
kernel_stats[kernel_name] += (record.end_timestamp - record.start_timestamp);
}
@@ -168,8 +168,8 @@ generate_stats(const output_config& /*cfg*/,
stats_entry_t
generate_stats(const output_config& /*cfg*/,
const metadata& tool_metadata,
const generator<rocprofiler_buffer_tracing_memory_allocation_record_t>& data)
const metadata& tool_metadata,
const generator<tool_buffer_tracing_memory_allocation_ext_record_t>& data)
{
auto memory_allocation_stats = stats_map_t{};
for(auto ditr : data)
projects/rocprofiler-sdk/source/lib/output/generateStats.hpp
+3 -3
Просмотреть файл
@@ -72,9 +72,9 @@ generate_stats(const output_config& cf
const generator<rocprofiler_buffer_tracing_rccl_api_record_t>& data);
stats_entry_t
generate_stats(const output_config& cfg,
const metadata& tool_metadata,
const generator<rocprofiler_buffer_tracing_memory_allocation_record_t>& data);
generate_stats(const output_config& cfg,
const metadata& tool_metadata,
const generator<tool_buffer_tracing_memory_allocation_ext_record_t>& data);
stats_entry_t
generate_stats(const output_config& cfg,
projects/rocprofiler-sdk/source/lib/output/generator.hpp
+106 -24
Просмотреть файл
@@ -25,12 +25,14 @@
#include "tmp_file_buffer.hpp"
#include "lib/common/container/ring_buffer.hpp"
#include "lib/common/demangle.hpp"
#include "lib/common/logging.hpp"
#include <fmt/format.h>
#include <iosfwd>
#include <mutex>
#include <numeric>
#include <set>
#include <vector>
@@ -76,35 +78,78 @@ get_buffer_elements(common::container::ring_buffer<Tp>&& buf)
template <typename Tp, domain_type DomainT>
struct buffered_output;
struct defer_size
{};
template <typename Tp>
struct generator
{
generator() = delete;
virtual ~generator() = default;
generator(const generator&) = delete;
generator(generator&&) noexcept = delete;
generator& operator=(const generator&) = delete;
generator& operator=(generator&&) noexcept = delete;
auto begin() { return m_pos.begin(); }
auto begin() const { return m_pos.begin(); }
auto cbegin() const { return m_pos.cbegin(); }
auto end() { return m_pos.end(); }
auto end() const { return m_pos.end(); }
auto cend() const { return m_pos.cend(); }
auto size() const { return m_pos.size(); }
auto empty() const { return m_pos.empty(); }
virtual std::vector<Tp> get(size_t) const = 0;
protected:
explicit generator(defer_size);
explicit generator(size_t sz);
void resize(size_t sz);
std::vector<size_t> m_pos = {};
};
template <typename Tp>
generator<Tp>::generator(defer_size)
{}
template <typename Tp>
generator<Tp>::generator(size_t sz)
{
resize(sz);
}
template <typename Tp>
void
generator<Tp>::resize(size_t sz)
{
m_pos.resize(sz, 0);
std::iota(m_pos.begin(), m_pos.end(), 0);
}
template <typename Tp>
struct file_generator : public generator<Tp>
{
template <typename Up, domain_type DomainT>
friend struct buffered_output;
generator() = delete;
~generator() = default;
file_generator() = delete;
~file_generator() override = default;
generator(const generator&) = delete;
generator(generator&&) = delete;
generator& operator=(const generator&) = delete;
generator& operator=(generator&&) = delete;
file_generator(const file_generator&) = delete;
file_generator(file_generator&&) noexcept = delete;
file_generator& operator=(const file_generator&) = delete;
file_generator& operator=(file_generator&&) noexcept = delete;
auto begin() { return file_pos.begin(); }
auto begin() const { return file_pos.begin(); }
auto cbegin() const { return file_pos.cbegin(); }
auto end() { return file_pos.end(); }
auto end() const { return file_pos.end(); }
auto cend() const { return file_pos.cend(); }
auto size() const { return file_pos.size(); }
auto empty() const { return file_pos.empty(); }
std::vector<Tp> get(std::streampos itr) const;
std::vector<Tp> get(size_t itr) const override;
private:
generator(file_buffer<Tp>* fbuf);
explicit file_generator(file_buffer<Tp>* fbuf);
file_buffer<Tp>* filebuf = nullptr;
std::lock_guard<std::mutex> lk_guard;
@@ -112,19 +157,38 @@ private:
};
template <typename Tp>
generator<Tp>::generator(file_buffer<Tp>* fbuf)
: filebuf{fbuf}
file_generator<Tp>::file_generator(file_buffer<Tp>* fbuf)
: generator<Tp>{fbuf->file.file_pos.size()}
, filebuf{fbuf}
, lk_guard{filebuf->file.file_mutex}
, file_pos{filebuf->file.file_pos}
{}
template <typename Tp>
std::vector<Tp>
generator<Tp>::get(std::streampos itr) const
file_generator<Tp>::get(size_t idx) const
{
auto _data = std::vector<Tp>{};
auto& _fs = filebuf->file.stream;
_fs.seekg(itr); // set to the absolute position
if(idx >= file_pos.size())
{
ROCP_ERROR << fmt::format("file_generator has no file position at index {}", idx);
return _data;
}
auto itr = file_pos.begin();
std::advance(itr, idx);
if(itr == file_pos.end())
{
ROCP_ERROR << fmt::format("file_generator at index {} == end of file_pos set", idx);
return _data;
}
ROCP_TRACE << fmt::format("file_generator file position at index={} :: ", idx) << *itr;
_fs.seekg(*itr); // set to the absolute position
if(!_fs.eof())
{
auto _buffer = ring_buffer_t<Tp>{};
@@ -138,9 +202,11 @@ generator<Tp>::get(std::streampos itr) const
namespace cereal
{
namespace tool = ::rocprofiler::tool;
template <typename ArchiveT, typename Tp>
void
save(ArchiveT& ar, const rocprofiler::tool::generator<Tp>& data)
save(ArchiveT& ar, const tool::file_generator<Tp>& data)
{
ar.makeArray();
for(auto itr : data)
@@ -150,4 +216,20 @@ save(ArchiveT& ar, const rocprofiler::tool::generator<Tp>& data)
ar(ditr);
}
}
template <typename ArchiveT, typename Tp>
void
save(ArchiveT& ar, const tool::generator<Tp>& data)
{
if(const auto* fdata = dynamic_cast<const tool::file_generator<Tp>*>(&data); fdata != nullptr)
{
save(ar, *fdata);
}
else
{
ROCP_CI_LOG(WARNING) << fmt::format(
"dynamic_cast failed for {}",
rocprofiler::common::cxx_demangle(typeid(tool::generator<Tp>).name()));
}
}
} // namespace cereal
projects/rocprofiler-sdk/source/lib/output/kernel_symbol_info.hpp
+17 -2
Просмотреть файл
@@ -82,17 +82,32 @@ using kernel_symbol_data_map_t = std::unordered_map<rocprofiler_kernel_id_t, ker
namespace cereal
{
#define SAVE_DATA_FIELD(FIELD) ar(make_nvp(#FIELD, data.FIELD))
#define LOAD_DATA_FIELD(FIELD) ar(make_nvp(#FIELD, data.FIELD))
template <typename ArchiveT>
void
save(ArchiveT& ar, const ::rocprofiler::tool::kernel_symbol_info& data)
{
cereal::save(ar,
static_cast<const ::rocprofiler::tool::rocprofiler_kernel_symbol_info_t&>(data));
using base_type = ::rocprofiler::tool::kernel_symbol_info::base_type;
cereal::save(ar, static_cast<const base_type&>(data));
SAVE_DATA_FIELD(formatted_kernel_name);
SAVE_DATA_FIELD(demangled_kernel_name);
SAVE_DATA_FIELD(truncated_kernel_name);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, ::rocprofiler::tool::kernel_symbol_info& data)
{
using base_type = ::rocprofiler::tool::kernel_symbol_info::base_type;
cereal::load(ar, static_cast<base_type&>(data));
LOAD_DATA_FIELD(formatted_kernel_name);
LOAD_DATA_FIELD(demangled_kernel_name);
LOAD_DATA_FIELD(truncated_kernel_name);
}
#undef SAVE_DATA_FIELD
#undef LOAD_DATA_FIELD
} // namespace cereal
projects/rocprofiler-sdk/source/lib/output/metadata.cpp
+33 -7
Просмотреть файл
@@ -21,14 +21,21 @@
// SOFTWARE.
#include "metadata.hpp"
#include "agent_info.hpp"
#include "host_symbol_info.hpp"
#include "kernel_symbol_info.hpp"
#include "node_info.hpp"
#include "lib/att-tool/att_lib_wrapper.hpp"
#include "lib/common/environment.hpp"
#include "lib/common/filesystem.hpp"
#include "lib/common/logging.hpp"
#include "lib/common/string_entry.hpp"
#include "lib/common/utility.hpp"
#include "lib/output/agent_info.hpp"
#include "lib/output/host_symbol_info.hpp"
#include "lib/output/kernel_symbol_info.hpp"
#include "lib/output/node_info.hpp"
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/rocprofiler.h>
@@ -37,6 +44,7 @@
#include <fmt/core.h>
#include <fmt/format.h>
#include <dlfcn.h>
#include <unistd.h>
#include <cstdint>
#include <memory>
@@ -46,9 +54,10 @@ namespace rocprofiler
{
namespace tool
{
namespace fs = common::filesystem;
namespace
{
namespace fs = ::rocprofiler::common::filesystem;
rocprofiler_status_t
query_pc_sampling_configuration(const rocprofiler_pc_sampling_configuration_t* configs,
long unsigned int num_config,
@@ -90,6 +99,8 @@ host_function_info::host_function_info()
metadata::metadata(inprocess)
: buffer_names{sdk::get_buffer_tracing_names()}
, callback_names{sdk::get_callback_tracing_names()}
, node_data{read_node_info()}
, command_line{common::read_command_line(getpid())}
{
ROCPROFILER_CHECK(rocprofiler_query_available_agents(
ROCPROFILER_AGENT_INFO_VERSION_0,
@@ -166,7 +177,7 @@ void metadata::init(inprocess)
{
auto _info = rocprofiler_counter_info_v1_t{};
auto _dim_ids = std::vector<rocprofiler_counter_dimension_id_t>{};
auto _dim_info = std::vector<rocprofiler_record_dimension_info_t>{};
auto _dim_info = std::vector<rocprofiler_counter_record_dimension_info_t>{};
ROCPROFILER_CHECK(rocprofiler_query_counter_info(
counters[i],
@@ -399,12 +410,12 @@ metadata::get_counter_dimension_info() const
_ret.emplace_back(ditr);
}
auto _sorter = [](const rocprofiler_record_dimension_info_t& lhs,
const rocprofiler_record_dimension_info_t& rhs) {
auto _sorter = [](const rocprofiler_counter_record_dimension_info_t& lhs,
const rocprofiler_counter_record_dimension_info_t& rhs) {
return std::tie(lhs.id, lhs.instance_size) < std::tie(rhs.id, rhs.instance_size);
};
auto _equiv = [](const rocprofiler_record_dimension_info_t& lhs,
const rocprofiler_record_dimension_info_t& rhs) {
auto _equiv = [](const rocprofiler_counter_record_dimension_info_t& lhs,
const rocprofiler_counter_record_dimension_info_t& rhs) {
return std::tie(lhs.id, lhs.instance_size) == std::tie(rhs.id, rhs.instance_size);
};
@@ -495,12 +506,27 @@ metadata::is_runtime_initialized(rocprofiler_runtime_initialization_operation_t
runtime_op);
}
void
metadata::set_process_id(pid_t _pid, pid_t _ppid, const std::vector<std::string>& _command_line)
{
process_id = _pid;
parent_process_id = _ppid;
if(!_command_line.empty())
command_line = _command_line;
else if(_pid == getpid())
command_line = common::read_command_line(_pid);
if(auto _start_ns = common::get_process_start_time_ns(_pid); _start_ns > 0)
process_start_ns = _start_ns;
}
std::string_view
metadata::get_marker_message(uint64_t corr_id) const
{
return marker_messages.rlock(
[](const auto& _data, uint64_t _corr_id_v) -> std::string_view {
return _data.at(_corr_id_v);
return (_data.count(_corr_id_v) > 0) ? _data.at(_corr_id_v) : std::string_view{};
},
corr_id);
}
projects/rocprofiler-sdk/source/lib/output/metadata.hpp
+9 -1
Просмотреть файл
@@ -33,6 +33,7 @@
#include "lib/common/demangle.hpp"
#include "lib/common/logging.hpp"
#include "lib/common/synchronized.hpp"
#include "lib/output/node_info.hpp"
#include <rocprofiler-sdk/agent.h>
#include <rocprofiler-sdk/buffer_tracing.h>
@@ -78,7 +79,7 @@ namespace tool
using marker_message_map_t = std::unordered_map<uint64_t, std::string>;
using marker_message_ordered_map_t = std::map<uint64_t, std::string>;
using string_entry_map_t = std::unordered_map<size_t, std::unique_ptr<std::string>>;
using counter_dimension_vec_t = std::vector<rocprofiler_record_dimension_info_t>;
using counter_dimension_vec_t = std::vector<rocprofiler_counter_record_dimension_info_t>;
using external_corr_id_set_t = std::unordered_set<uint64_t>;
using code_obj_decoder_t = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate;
using instruction_t = rocprofiler::sdk::codeobj::disassembly::Instruction;
@@ -121,6 +122,7 @@ struct metadata
{};
pid_t process_id = 0;
pid_t parent_process_id = 0;
uint64_t process_start_ns = 0;
uint64_t process_end_ns = 0;
agent_info_vec_t agents = {};
@@ -140,6 +142,8 @@ struct metadata
att_filenames_map_t att_filenames = {};
synced_obj<pc_sampling_stats_t> pc_sampling_stats = {};
synced_obj<runtime_initialization_set_t> runtime_initialization_set = {};
node_info node_data = {};
std::vector<std::string> command_line = {};
metadata() = default;
metadata(inprocess);
@@ -187,6 +191,10 @@ struct metadata
bool add_external_correlation_id(uint64_t);
bool add_runtime_initialization(rocprofiler_runtime_initialization_operation_t);
void set_process_id(pid_t _pid,
pid_t _ppid = 0,
const std::vector<std::string>& _command_line = {});
std::string_view get_marker_message(uint64_t corr_id) const;
std::string_view get_kernel_name(uint64_t kernel_id, uint64_t rename_id) const;
std::string_view get_kind_name(rocprofiler_callback_tracing_kind_t kind) const;
projects/rocprofiler-sdk/source/lib/output/node_info.cpp
+79
Просмотреть файл
@@ -0,0 +1,79 @@
// 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/output/node_info.hpp"
#include "lib/common/logging.hpp"
#include <rocprofiler-sdk/cxx/details/tokenize.hpp>
#include <sys/utsname.h>
#include <fstream>
namespace rocprofiler
{
namespace tool
{
using utsname_t = struct utsname;
node_info&
read_node_info(node_info& _info)
{
{
if(auto ifs = std::ifstream{"/etc/machine-id"})
{
auto _mach_id = std::string{};
if((ifs >> _mach_id) && !_mach_id.empty())
_info.machine_id = sdk::parse::strip(std::move(_mach_id), "\n\t\r ");
}
}
auto _sys_info = utsname_t{};
if(uname(&_sys_info) == 0)
{
auto _assign = [](auto& _dst, const char* _src) {
if(_src) _dst = std::string{_src};
};
_assign(_info.system_name, _sys_info.sysname);
_assign(_info.hostname, _sys_info.nodename);
_assign(_info.release, _sys_info.release);
_assign(_info.version, _sys_info.version);
_assign(_info.hardware_name, _sys_info.machine);
_assign(_info.domain_name, _sys_info.domainname);
}
else
{
ROCP_WARNING << "error retrieving uname info";
}
return _info;
}
node_info
read_node_info()
{
auto _val = node_info{};
return read_node_info(_val);
}
} // namespace tool
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/output/node_info.hpp
+102
Просмотреть файл
@@ -0,0 +1,102 @@
// 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/agent.h>
#include <rocprofiler-sdk/cxx/serialization.hpp>
#include <unordered_map>
#include <vector>
namespace rocprofiler
{
namespace tool
{
struct node_info
{
node_info() = default;
~node_info() = default;
node_info(const node_info&) = default;
node_info(node_info&&) noexcept = default;
node_info& operator=(const node_info&) = default;
node_info& operator=(node_info&&) noexcept = default;
uint64_t id = 0;
uint64_t hash = 0;
std::string machine_id = {};
std::string system_name = {};
std::string hostname = {};
std::string release = {};
std::string version = {};
std::string hardware_name = {};
std::string domain_name = {};
};
node_info&
read_node_info(node_info& _info);
node_info
read_node_info();
using node_info_vec_t = std::vector<node_info>;
} // namespace tool
} // namespace rocprofiler
namespace cereal
{
#define SAVE_DATA_FIELD(FIELD) ar(make_nvp(#FIELD, data.FIELD))
#define LOAD_DATA_FIELD(FIELD) ar(make_nvp(#FIELD, data.FIELD))
template <typename ArchiveT>
void
save(ArchiveT& ar, const ::rocprofiler::tool::node_info& data)
{
SAVE_DATA_FIELD(id);
SAVE_DATA_FIELD(hash);
SAVE_DATA_FIELD(machine_id);
SAVE_DATA_FIELD(system_name);
SAVE_DATA_FIELD(hostname);
SAVE_DATA_FIELD(release);
SAVE_DATA_FIELD(version);
SAVE_DATA_FIELD(hardware_name);
SAVE_DATA_FIELD(domain_name);
}
template <typename ArchiveT>
void
load(ArchiveT& ar, ::rocprofiler::tool::node_info& data)
{
LOAD_DATA_FIELD(id);
LOAD_DATA_FIELD(hash);
LOAD_DATA_FIELD(machine_id);
LOAD_DATA_FIELD(system_name);
LOAD_DATA_FIELD(hostname);
LOAD_DATA_FIELD(release);
LOAD_DATA_FIELD(version);
LOAD_DATA_FIELD(hardware_name);
LOAD_DATA_FIELD(domain_name);
}
#undef SAVE_DATA_FIELD
#undef LOAD_DATA_FIELD
} // namespace cereal
projects/rocprofiler-sdk/source/lib/output/output_key.cpp
+26 -21
Просмотреть файл
@@ -223,22 +223,26 @@ output_keys(std::string _tag)
}
for(auto&& itr : std::initializer_list<output_key>{
{"%argv%", _argv_string, "Entire command-line condensed into a single string"},
{"%argt%",
{"argv", _argv_string, "Entire command-line condensed into a single string"},
{"argt",
_argt_string,
"Similar to `%argv%` except basename of first command line argument"},
{"%args%", _args_string, "All command line arguments condensed into a single string"},
{"%tag%", _tag0_string, "Basename of first command line argument"}})
{"args", _args_string, "All command line arguments condensed into a single string"},
{"tag", _tag0_string, "Basename of first command line argument"}})
{
_options.emplace_back(itr);
_options.emplace_back(fmt::format("%{}%", itr.key), itr.value, itr.description);
_options.emplace_back(fmt::format("{}{}{}", '{', itr.key, '}'), itr.value, itr.description);
}
if(!_cmdline.empty())
{
for(size_t i = 0; i < _cmdline.size(); ++i)
{
auto _v = _cmdline.at(i);
_options.emplace_back(fmt::format("%arg{}%", i), _v, fmt::format("Argument #{}", i));
auto _v = _cmdline.at(i);
auto itr = output_key{fmt::format("arg{}", i), _v, fmt::format("Argument #{}", i)};
_options.emplace_back(fmt::format("%{}%", itr.key), itr.value, itr.description);
_options.emplace_back(
fmt::format("{}{}{}", '{', itr.key, '}'), itr.value, itr.description);
}
}
@@ -247,22 +251,23 @@ output_keys(std::string _tag)
auto _hostname = get_hostname();
for(auto&& itr : std::initializer_list<output_key>{
{"%hostname%", _hostname, "Network hostname"},
{"%pid%", _proc_id, "Process identifier"},
{"%ppid%", _parent_id, "Parent process identifier"},
{"%pgid%", _pgroup_id, "Process group identifier"},
{"%psid%", _session_id, "Process session identifier"},
{"%psize%", _proc_size, "Number of sibling process"},
{"%job%", _slurm_job_id, "SLURM_JOB_ID env variable"},
{"%rank%", _slurm_proc_id, "MPI/UPC++ rank"},
{"%size%", _dmp_size, "MPI/UPC++ size"},
{"%nid%", _uniq_id, "%rank% if possible, otherwise %pid%"},
{"%cwd%", fs::current_path().string(), "Current working path"},
{"%launch_date%", _launch_date, "Date according to date format ROCPROF_DATE_FORMAT"},
{"%launch_time%", _launch_time, "Date and/or time according to ROCPROF_TIME_FORMAT"},
{"hostname", _hostname, "Network hostname"},
{"pid", _proc_id, "Process identifier"},
{"ppid", _parent_id, "Parent process identifier"},
{"pgid", _pgroup_id, "Process group identifier"},
{"psid", _session_id, "Process session identifier"},
{"psize", _proc_size, "Number of sibling process"},
{"job", _slurm_job_id, "SLURM_JOB_ID env variable"},
{"rank", _slurm_proc_id, "MPI/UPC++ rank"},
{"size", _dmp_size, "MPI/UPC++ size"},
{"nid", _uniq_id, "%rank% if possible, otherwise %pid%"},
{"cwd", fs::current_path().string(), "Current working path"},
{"launch_date", _launch_date, "Date according to date format ROCPROF_DATE_FORMAT"},
{"launch_time", _launch_time, "Date and/or time according to ROCPROF_TIME_FORMAT"},
})
{
_options.emplace_back(itr);
_options.emplace_back(fmt::format("%{}%", itr.key), itr.value, itr.description);
_options.emplace_back(fmt::format("{}{}{}", '{', itr.key, '}'), itr.value, itr.description);
}
for(auto&& itr : std::initializer_list<output_key>{
projects/rocprofiler-sdk/source/lib/output/output_stream.cpp
+5 -2
Просмотреть файл
@@ -87,7 +87,10 @@ get_output_filename(const output_config& cfg, std::string_view fname, std::strin
}
output_stream
get_output_stream(const output_config& cfg, std::string_view fname, std::string_view ext)
get_output_stream(const output_config& cfg,
std::string_view fname,
std::string_view ext,
std::ios::openmode mode)
{
auto cfg_output_path = tool::format_path(cfg.output_path);
@@ -99,7 +102,7 @@ get_output_stream(const output_config& cfg, std::string_view fname, std::string_
return {&std::clog, [](auto*&) {}};
auto output_file = get_output_filename(cfg, fname, ext);
auto* _ofs = new(std::nothrow) std::ofstream{output_file};
auto* _ofs = new(std::nothrow) std::ofstream{output_file, mode};
LOG_IF(FATAL, !_ofs) << fmt::format("Failed to allocate ofstream for output file '{}'",
output_file);
projects/rocprofiler-sdk/source/lib/output/output_stream.hpp
+4 -1
Просмотреть файл
@@ -80,6 +80,9 @@ std::string
get_output_filename(const output_config& cfg, std::string_view fname, std::string_view ext);
output_stream
get_output_stream(const output_config& cfg, std::string_view fname, std::string_view ext);
get_output_stream(const output_config& cfg,
std::string_view fname,
std::string_view ext,
std::ios::openmode mode = {});
} // namespace tool
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/output/stream_info.hpp
+40 -8
Просмотреть файл
@@ -41,14 +41,12 @@ struct tool_buffer_tracing_kernel_dispatch_ext_record_t
using base_type = rocprofiler_buffer_tracing_kernel_dispatch_record_t;
tool_buffer_tracing_kernel_dispatch_ext_record_t(const base_type& _base,
const rocprofiler_stream_id_t& _stream_id,
const uint64_t& _kernel_rename_val)
const rocprofiler_stream_id_t& _stream_id)
: base_type{_base}
, stream_id{_stream_id}
, kernel_rename_val{_kernel_rename_val}
{}
tool_buffer_tracing_kernel_dispatch_ext_record_t();
tool_buffer_tracing_kernel_dispatch_ext_record_t() = delete;
~tool_buffer_tracing_kernel_dispatch_ext_record_t() = default;
tool_buffer_tracing_kernel_dispatch_ext_record_t(
const tool_buffer_tracing_kernel_dispatch_ext_record_t&) = default;
@@ -59,8 +57,7 @@ struct tool_buffer_tracing_kernel_dispatch_ext_record_t
tool_buffer_tracing_kernel_dispatch_ext_record_t& operator =(
tool_buffer_tracing_kernel_dispatch_ext_record_t&&) noexcept = default;
rocprofiler_stream_id_t stream_id = {};
uint64_t kernel_rename_val = {};
rocprofiler_stream_id_t stream_id = {.handle = 0};
};
struct tool_buffer_tracing_memory_copy_ext_record_t
@@ -74,7 +71,7 @@ struct tool_buffer_tracing_memory_copy_ext_record_t
, stream_id{_stream_id}
{}
tool_buffer_tracing_memory_copy_ext_record_t();
tool_buffer_tracing_memory_copy_ext_record_t() = delete;
~tool_buffer_tracing_memory_copy_ext_record_t() = default;
tool_buffer_tracing_memory_copy_ext_record_t(
const tool_buffer_tracing_memory_copy_ext_record_t&) = default;
@@ -87,6 +84,32 @@ struct tool_buffer_tracing_memory_copy_ext_record_t
rocprofiler_stream_id_t stream_id = {};
};
struct tool_buffer_tracing_memory_allocation_ext_record_t
: rocprofiler_buffer_tracing_memory_allocation_record_t
{
using base_type = rocprofiler_buffer_tracing_memory_allocation_record_t;
tool_buffer_tracing_memory_allocation_ext_record_t(const base_type& _base,
const rocprofiler_stream_id_t& _stream_id)
: base_type{_base}
, stream_id{_stream_id}
{}
tool_buffer_tracing_memory_allocation_ext_record_t() = delete;
~tool_buffer_tracing_memory_allocation_ext_record_t() = default;
tool_buffer_tracing_memory_allocation_ext_record_t(
const tool_buffer_tracing_memory_allocation_ext_record_t&) = default;
tool_buffer_tracing_memory_allocation_ext_record_t(
tool_buffer_tracing_memory_allocation_ext_record_t&&) noexcept = default;
tool_buffer_tracing_memory_allocation_ext_record_t& operator =(
const tool_buffer_tracing_memory_allocation_ext_record_t&) = default;
tool_buffer_tracing_memory_allocation_ext_record_t& operator =(
tool_buffer_tracing_memory_allocation_ext_record_t&&) noexcept = default;
rocprofiler_stream_id_t stream_id = {};
};
} // namespace tool
} // namespace rocprofiler
@@ -101,7 +124,6 @@ save(ArchiveT&
{
cereal::save(ar, static_cast<const rocprofiler_buffer_tracing_kernel_dispatch_record_t&>(data));
SAVE_DATA_FIELD(stream_id);
SAVE_DATA_FIELD(kernel_rename_val);
}
template <typename ArchiveT>
@@ -112,5 +134,15 @@ save(ArchiveT& ar, const ::rocprofiler::tool::tool_buffer_tracing_memory_copy_ex
SAVE_DATA_FIELD(stream_id);
}
template <typename ArchiveT>
void
save(ArchiveT& ar,
const ::rocprofiler::tool::tool_buffer_tracing_memory_allocation_ext_record_t& data)
{
cereal::save(ar,
static_cast<const rocprofiler_buffer_tracing_memory_allocation_record_t&>(data));
SAVE_DATA_FIELD(stream_id);
}
#undef SAVE_DATA_FIELD
} // namespace cereal
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/stream_stack.cpp
-2
Просмотреть файл
@@ -42,7 +42,6 @@ get_stream_stack()
construct(rocprofiler_stream_id_t{.handle = 0});
return _v;
}
} // namespace
void
@@ -69,7 +68,6 @@ stream_stack_empty()
{
return CHECK_NOTNULL(get_stream_stack())->empty();
}
} // namespace stream
} // namespace tool
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp
+180 -114
Просмотреть файл
@@ -20,8 +20,6 @@
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
// SOFTWARE.
#include "rocprofiler-sdk/defines.h"
#include "rocprofiler-sdk/dispatch_counting_service.h"
#define _GNU_SOURCE 1
#define _DEFAULT_SOURCE 1
@@ -35,7 +33,9 @@
#include "lib/common/filesystem.hpp"
#include "lib/common/logging.hpp"
#include "lib/common/scope_destructor.hpp"
#include "lib/common/simple_timer.hpp"
#include "lib/common/static_object.hpp"
#include "lib/common/static_tl_object.hpp"
#include "lib/common/string_entry.hpp"
#include "lib/common/synchronized.hpp"
#include "lib/common/units.hpp"
@@ -50,6 +50,7 @@
#include "lib/output/generateOTF2.hpp"
#include "lib/output/generatePerfetto.hpp"
#include "lib/output/generateStats.hpp"
#include "lib/output/metadata.hpp"
#include "lib/output/output_stream.hpp"
#include "lib/output/statistics.hpp"
#include "lib/output/stream_info.hpp"
@@ -60,10 +61,13 @@
#include <rocprofiler-sdk/agent.h>
#include <rocprofiler-sdk/buffer_tracing.h>
#include <rocprofiler-sdk/callback_tracing.h>
#include <rocprofiler-sdk/defines.h>
#include <rocprofiler-sdk/dispatch_counting_service.h>
#include <rocprofiler-sdk/experimental/counters.h>
#include <rocprofiler-sdk/experimental/thread_trace.h>
#include <rocprofiler-sdk/external_correlation.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/intercept_table.h>
#include <rocprofiler-sdk/internal_threading.h>
#include <rocprofiler-sdk/marker/api_id.h>
#include <rocprofiler-sdk/rocprofiler.h>
@@ -79,6 +83,7 @@
#include <cassert>
#include <chrono>
#include <csignal>
#include <cstdlib>
#include <cstring>
#include <fstream>
#include <future>
@@ -261,11 +266,11 @@ thread_local auto thread_dispatch_rename_dtor = common::scope_destructor{[]() {
thread_dispatch_rename = nullptr;
}};
// Stores stream_ids and kernel_rename_vals for kernel-rename service and hip stream display service
struct kernel_rename_and_stream_display_pair
// Stores stream ids and kernel region ids for kernel-rename service and hip stream display service
struct kernel_rename_and_stream_data
{
uint64_t kernel_rename_val{0};
rocprofiler_stream_id_t stream_id{.handle = 0};
uint64_t region_id = 0; // roctx region correlation id
rocprofiler_stream_id_t stream_id = {.handle = 0};
};
bool
@@ -451,58 +456,60 @@ record_execution_profile(rocprofiler_thread_id_t thr_
return 0;
}
int
set_kernel_rename_and_stream_display_correlation_id(
rocprofiler_thread_id_t thr_id,
rocprofiler_context_id_t ctx_id,
rocprofiler_external_correlation_id_request_kind_t kind,
rocprofiler_tracing_operation_t op,
uint64_t internal_corr_id,
rocprofiler_user_data_t* external_corr_id,
void* user_data)
template <typename Tp>
rocprofiler_stream_id_t
get_stream_id(Tp* _record)
{
auto _stream_id = rocprofiler_stream_id_t{.handle = 0};
if(_record->correlation_id.external.ptr != nullptr)
{
// Extract the stream id
auto* _ecid_data =
static_cast<kernel_rename_and_stream_data*>(_record->correlation_id.external.ptr);
_stream_id = _ecid_data->stream_id;
auto _region_id = _ecid_data->region_id;
_record->correlation_id.external.value = _region_id;
delete _ecid_data;
}
return _stream_id;
}
int
set_kernel_rename_and_stream_correlation_id(rocprofiler_thread_id_t thr_id,
rocprofiler_context_id_t ctx_id,
rocprofiler_external_correlation_id_request_kind_t kind,
rocprofiler_tracing_operation_t op,
uint64_t internal_corr_id,
rocprofiler_user_data_t* external_corr_id,
void* user_data)
{
ROCP_FATAL_IF(kind != ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_KERNEL_DISPATCH &&
kind != ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MEMORY_COPY)
<< "unexpected kind: " << kind;
// Check whether services are enabled
const bool kernel_rename_service_enabled =
kind == ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_KERNEL_DISPATCH &&
tool::get_config().kernel_rename && thread_dispatch_rename != nullptr &&
!thread_dispatch_rename->empty() &&
kind == ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_KERNEL_DISPATCH;
const bool hip_stream_display_enabled = !tool::get_config().group_by_queue;
if(kernel_rename_service_enabled && ctx_id == counter_collection_ctx)
!thread_dispatch_rename->empty();
const bool hip_stream_enabled = !tool::get_config().group_by_queue;
if(!kernel_rename_service_enabled && !hip_stream_enabled) return 1;
auto* _info = new kernel_rename_and_stream_data{};
// Get value for kernel rename service
if(kernel_rename_service_enabled)
{
auto val = thread_dispatch_rename->top();
if(tool_metadata) tool_metadata->add_external_correlation_id(val);
external_corr_id->value = val;
return 0;
}
// End early if counter collection context and no kernel rename service is enabled
if(ctx_id == counter_collection_ctx)
{
return 0;
_info->region_id = thread_dispatch_rename->top();
if(tool_metadata) tool_metadata->add_external_correlation_id(_info->region_id);
}
kernel_rename_and_stream_display_pair* kernel_rename_and_stream_display_vals = nullptr;
if(kernel_rename_service_enabled || hip_stream_display_enabled)
{
kernel_rename_and_stream_display_vals = new kernel_rename_and_stream_display_pair{};
}
// Get value for kernel rename service
if(kernel_rename_service_enabled && kernel_rename_and_stream_display_vals != nullptr)
{
auto val = thread_dispatch_rename->top();
if(tool_metadata) tool_metadata->add_external_correlation_id(val);
kernel_rename_and_stream_display_vals->kernel_rename_val = val;
}
// Get stream ID from stream HIP display service
if(hip_stream_display_enabled && kernel_rename_and_stream_display_vals != nullptr)
if(hip_stream_enabled)
{
auto stream_id = tool::stream::get_stream_id();
kernel_rename_and_stream_display_vals->stream_id = stream_id;
_info->stream_id = rocprofiler::tool::stream::get_stream_id();
}
// Set the external correlation id service to point to struct
external_corr_id->ptr = kernel_rename_and_stream_display_vals;
external_corr_id->ptr = _info;
common::consume_args(thr_id, ctx_id, kind, op, internal_corr_id, user_data);
@@ -1005,25 +1012,12 @@ buffered_tracing_callback(rocprofiler_context_id_t /*context*/,
{
auto* record = static_cast<rocprofiler_buffer_tracing_kernel_dispatch_record_t*>(
header->payload);
rocprofiler_stream_id_t stream_id{.handle = 0};
uint64_t kernel_rename_val = 0;
if((!tool::get_config().group_by_queue || tool::get_config().kernel_rename) &&
record->correlation_id.external.ptr != nullptr)
{
// Extract the stream id
auto* kernel_stream_pair_ptr =
static_cast<kernel_rename_and_stream_display_pair*>(
record->correlation_id.external.ptr);
stream_id = kernel_stream_pair_ptr->stream_id;
kernel_rename_val = kernel_stream_pair_ptr->kernel_rename_val;
delete kernel_stream_pair_ptr;
record->correlation_id.external.value = kernel_rename_val;
}
rocprofiler::tool::tool_buffer_tracing_kernel_dispatch_ext_record_t
record_with_stream{*record, stream_id, kernel_rename_val};
tool::write_ring_buffer(record_with_stream, domain_type::KERNEL_DISPATCH);
}
auto stream_id = get_stream_id(record);
tool::write_ring_buffer(
tool::tool_buffer_tracing_kernel_dispatch_ext_record_t{*record, stream_id},
domain_type::KERNEL_DISPATCH);
}
else if(header->kind == ROCPROFILER_BUFFER_TRACING_HSA_CORE_API ||
header->kind == ROCPROFILER_BUFFER_TRACING_HSA_AMD_EXT_API ||
header->kind == ROCPROFILER_BUFFER_TRACING_HSA_IMAGE_EXT_API ||
@@ -1038,28 +1032,21 @@ buffered_tracing_callback(rocprofiler_context_id_t /*context*/,
{
auto* record =
static_cast<rocprofiler_buffer_tracing_memory_copy_record_t*>(header->payload);
rocprofiler_stream_id_t stream_id{.handle = 0};
if(!tool::get_config().group_by_queue &&
record->correlation_id.external.ptr != nullptr)
{
// Extract the stream id
auto* kernel_stream_pair_ptr =
static_cast<kernel_rename_and_stream_display_pair*>(
record->correlation_id.external.ptr);
stream_id = kernel_stream_pair_ptr->stream_id;
delete kernel_stream_pair_ptr;
record->correlation_id.external.ptr = nullptr;
}
rocprofiler::tool::tool_buffer_tracing_memory_copy_ext_record_t record_with_stream{
*record, stream_id};
tool::write_ring_buffer(record_with_stream, domain_type::MEMORY_COPY);
auto stream_id = get_stream_id(record);
tool::write_ring_buffer(
tool::tool_buffer_tracing_memory_copy_ext_record_t{*record, stream_id},
domain_type::MEMORY_COPY);
}
else if(header->kind == ROCPROFILER_BUFFER_TRACING_MEMORY_ALLOCATION)
{
auto* record = static_cast<rocprofiler_buffer_tracing_memory_allocation_record_t*>(
header->payload);
tool::write_ring_buffer(*record, domain_type::MEMORY_ALLOCATION);
auto stream_id = get_stream_id(record);
tool::write_ring_buffer(
tool::tool_buffer_tracing_memory_allocation_ext_record_t{*record, stream_id},
domain_type::MEMORY_ALLOCATION);
}
else if(header->kind == ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY)
{
@@ -1444,11 +1431,12 @@ counter_record_callback(rocprofiler_dispatch_counting_service_data_t dispatch_da
auto counter_record = tool::tool_counter_record_t{};
counter_record.dispatch_data = dispatch_data;
counter_record.thread_id = user_data.value;
counter_record.kernel_rename_val = dispatch_data.correlation_id.external.value;
counter_record.dispatch_data.correlation_id.external.value = counter_record.kernel_rename_val;
auto serialized_records = std::vector<tool::tool_counter_value_t>{};
// must call get_stream_id on dispatch_data before copying to counter_record.dispatch_data
// so that external correlation id is updated before copy is made
counter_record.stream_id = get_stream_id(&dispatch_data);
counter_record.dispatch_data = dispatch_data;
counter_record.thread_id = user_data.value;
auto serialized_records = std::vector<tool::tool_counter_value_t>{};
serialized_records.reserve(record_count);
for(size_t count = 0; count < record_count; ++count)
@@ -1686,6 +1674,8 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
static constexpr auto null_context_id = rocprofiler_context_id_t{.handle = 0};
static constexpr auto null_buffer_id = rocprofiler_buffer_id_t{.handle = 0};
auto _init_timer = common::simple_timer{"[rocprofv3] tool initialization"};
client_finalizer = fini_func;
const uint64_t buffer_size = 16 * common::units::get_page_size();
@@ -2027,15 +2017,16 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
tool::get_config().benchmark_mode != tool::config::benchmark::execution_profile)
{
auto external_corr_id_request_kinds =
std::array<rocprofiler_external_correlation_id_request_kind_t, 2>{
std::array<rocprofiler_external_correlation_id_request_kind_t, 3>{
ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_KERNEL_DISPATCH,
ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MEMORY_COPY};
ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MEMORY_COPY,
ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MEMORY_ALLOCATION};
ROCPROFILER_CALL(rocprofiler_configure_external_correlation_id_request_service(
get_client_ctx(),
external_corr_id_request_kinds.data(),
external_corr_id_request_kinds.size(),
set_kernel_rename_and_stream_display_correlation_id,
set_kernel_rename_and_stream_correlation_id,
nullptr),
"Could not configure external correlation id request service");
@@ -2049,7 +2040,7 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
counter_collection_ctx,
counter_external_corr_id_request_kinds.data(),
counter_external_corr_id_request_kinds.size(),
set_kernel_rename_and_stream_display_correlation_id,
set_kernel_rename_and_stream_correlation_id,
nullptr),
"Could not configure external correlation id request service");
}
@@ -2115,12 +2106,47 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
start_context(get_client_ctx(), "primary rocprofv3");
}
tool_metadata->process_id = getpid();
rocprofiler_get_timestamp(&(tool_metadata->process_start_ns));
tool_metadata->set_process_id(getpid(), getppid());
// set_process_id should set process_start_ns unless it cannot read from /proc/<pid>/stat
if(tool_metadata->process_start_ns == 0)
rocprofiler_get_timestamp(&(tool_metadata->process_start_ns));
return 0;
}
void
api_timestamps_callback(rocprofiler_intercept_table_t table_id,
uint64_t lib_version,
uint64_t lib_instance,
void** /*tables*/,
uint64_t /*num_tables*/,
void* /*user_data*/)
{
static auto _once = std::once_flag{};
// compute major/minor/patch version info
uint32_t major = lib_version / 10000;
uint32_t minor = (lib_version % 10000) / 100;
uint32_t patch = lib_version % 100;
const char* table_name = nullptr;
ROCPROFILER_CHECK(rocprofiler_query_intercept_table_name(table_id, &table_name, nullptr));
ROCP_WARNING_IF(table_id != ROCPROFILER_MARKER_CONTROL_TABLE &&
table_id != ROCPROFILER_MARKER_NAME_TABLE && table_name)
<< fmt::format("{} version {}.{}.{} initialized (instance={})",
table_name,
major,
minor,
patch,
lib_instance);
std::call_once(_once, []() {
if(CHECK_NOTNULL(tool_metadata)->process_start_ns == 0)
rocprofiler_get_timestamp(&(tool_metadata->process_start_ns));
});
}
using stats_data_t = tool::stats_data_t;
using stats_entry_t = tool::stats_entry_t;
using domain_stats_vec_t = tool::domain_stats_vec_t;
@@ -2261,26 +2287,33 @@ generate_output(tool::buffered_output<Tp, DomainT>& output_v,
void
tool_fini(void* /*tool_data*/)
{
static bool _first = true;
if(!_first) return;
_first = false;
client_identifier = nullptr;
client_finalizer = nullptr;
tool_metadata->process_id = getpid();
rocprofiler_get_timestamp(&(tool_metadata->process_end_ns));
auto _fini_timer = common::simple_timer{"[rocprofv3] tool finalization"};
if(tool_metadata->process_end_ns == 0)
rocprofiler_get_timestamp(&(tool_metadata->process_end_ns));
flush();
rocprofiler_stop_context(get_client_ctx());
flush();
auto kernel_dispatch_output = rocprofiler::tool::kernel_dispatch_buffered_output_with_stream_t{
tool::get_config().kernel_trace};
auto kernel_dispatch_output =
rocprofiler::tool::kernel_dispatch_buffered_output_ext_t{tool::get_config().kernel_trace};
auto hsa_output = tool::hsa_buffered_output_t{tool::get_config().hsa_core_api_trace ||
tool::get_config().hsa_amd_ext_api_trace ||
tool::get_config().hsa_image_ext_api_trace ||
tool::get_config().hsa_finalizer_ext_api_trace};
auto hip_output = tool::hip_buffered_output_t{tool::get_config().hip_runtime_api_trace ||
tool::get_config().hip_compiler_api_trace};
auto memory_copy_output_output =
tool::memory_copy_buffered_output_with_stream_t{tool::get_config().memory_copy_trace};
auto memory_copy_output =
tool::memory_copy_buffered_output_ext_t{tool::get_config().memory_copy_trace};
auto marker_output = tool::marker_buffered_output_t{tool::get_config().marker_api_trace};
auto counters_output =
tool::counter_collection_buffered_output_t{tool::get_config().counter_collection};
@@ -2326,7 +2359,7 @@ tool_fini(void* /*tool_data*/)
generate_output(kernel_dispatch_output, outdata, contributions, cleanups);
generate_output(hsa_output, outdata, contributions, cleanups);
generate_output(hip_output, outdata, contributions, cleanups);
generate_output(memory_copy_output_output, outdata, contributions, cleanups);
generate_output(memory_copy_output, outdata, contributions, cleanups);
generate_output(memory_allocation_output, outdata, contributions, cleanups);
generate_output(marker_output, outdata, contributions, cleanups);
generate_output(rccl_output, outdata, contributions, cleanups);
@@ -2372,7 +2405,7 @@ tool_fini(void* /*tool_data*/)
hip_output.get_generator(),
hsa_output.get_generator(),
kernel_dispatch_output.get_generator(),
memory_copy_output_output.get_generator(),
memory_copy_output.get_generator(),
counters_output.get_generator(),
marker_output.get_generator(),
scratch_memory_output.get_generator(),
@@ -2396,7 +2429,7 @@ tool_fini(void* /*tool_data*/)
hip_output.get_generator(),
hsa_output.get_generator(),
kernel_dispatch_output.get_generator(),
memory_copy_output_output.get_generator(),
memory_copy_output.get_generator(),
counters_output.get_generator(),
marker_output.get_generator(),
scratch_memory_output.get_generator(),
@@ -2412,7 +2445,7 @@ tool_fini(void* /*tool_data*/)
auto hip_elem_data = hip_output.load_all();
auto hsa_elem_data = hsa_output.load_all();
auto kernel_dispatch_elem_data = kernel_dispatch_output.load_all();
auto memory_copy_elem_data = memory_copy_output_output.load_all();
auto memory_copy_elem_data = memory_copy_output.load_all();
auto marker_elem_data = marker_output.load_all();
auto scratch_memory_elem_data = scratch_memory_output.load_all();
auto rccl_elem_data = rccl_output.load_all();
@@ -2766,17 +2799,17 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext)
if(get_chained_signals().at(signo))
{
ROCP_INFO << fmt::format(
"[PPID={}][PID={}][TID={}][{}] rocprofv3 found chained signal for {}",
"[PPID={}][PID={}][TID={}][{}] rocprofv3 found chained signal handler for {}",
this_ppid,
this_pid,
this_tid,
this_func,
signo);
auto& _chained = *get_chained_signals().at(signo);
if(_chained.action)
if(auto& _chained = *get_chained_signals().at(signo); _chained.action)
{
ROCP_TRACE << fmt::format("[PPID={}][PID={}][TID={}][{}] rocprofv3 found chained "
"signal for {}... executing chained sigaction",
"signal handler for {}... executing chained sigaction",
this_ppid,
this_pid,
this_tid,
@@ -2786,8 +2819,8 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext)
_chained.action->sa_sigaction)
{
ROCP_TRACE << fmt::format(
"[PPID={}][PID={}][TID={}][{}] rocprofv3 found chained "
"signal for {}... executing chained sigaction (SIGINFO)",
"[PPID={}][PID={}][TID={}][{}] rocprofv3 found chained signal handler for "
"{}... executing chained sigaction (SIGINFO)",
this_ppid,
this_pid,
this_tid,
@@ -2796,11 +2829,12 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext)
_chained.action->sa_sigaction(signo, info, ucontext);
}
else if((_chained.action->sa_flags & SA_SIGINFO) != SA_SIGINFO &&
_chained.action->sa_handler)
_chained.action->sa_handler &&
_chained.action->sa_sigaction != &rocprofv3_error_signal_handler)
{
ROCP_TRACE << fmt::format(
"[PPID={}][PID={}][TID={}][{}] rocprofv3 found chained "
"signal for {}... executing chained sigaction (HANDLER)",
"[PPID={}][PID={}][TID={}][{}] rocprofv3 found chained signal handler for "
"{}... executing chained sigaction (HANDLER)",
this_ppid,
this_pid,
this_tid,
@@ -2814,8 +2848,8 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext)
if(_chained.handler)
{
ROCP_TRACE << fmt::format(
"[PPID={}][PID={}][TID={}][{}] rocprofv3 found chained "
"signal for {}... executing chained handler",
"[PPID={}][PID={}][TID={}][{}] rocprofv3 found chained signal handler for "
"{}... executing chained handler",
this_ppid,
this_pid,
this_tid,
@@ -2886,6 +2920,18 @@ rocprofiler_configure(uint32_t version,
"Loading extra counters");
}
int libs = ROCPROFILER_HSA_TABLE;
if(tool::get_config().hip_compiler_api_trace) libs |= ROCPROFILER_HIP_COMPILER_TABLE;
if(tool::get_config().hip_runtime_api_trace) libs |= ROCPROFILER_HIP_RUNTIME_TABLE;
if(tool::get_config().rccl_api_trace) libs |= ROCPROFILER_RCCL_TABLE;
if(tool::get_config().marker_api_trace) libs |= ROCPROFILER_MARKER_CORE_TABLE;
if(tool::get_config().rocdecode_api_trace) libs |= ROCPROFILER_ROCDECODE_TABLE;
if(tool::get_config().rocjpeg_api_trace) libs |= ROCPROFILER_ROCJPEG_TABLE;
ROCPROFILER_CALL(
rocprofiler_at_intercept_table_registration(api_timestamps_callback, libs, nullptr),
"api registration");
ROCP_INFO << id->name << " is using rocprofiler-sdk v" << major << "." << minor << "." << patch
<< " (" << runtime_version << ")";
@@ -2943,7 +2989,10 @@ rocprofv3_sigaction(int signum,
if(!is_handled_signal(signum) || !act || !tool::get_config().enable_signal_handlers)
return CHECK_NOTNULL(get_sigaction_function())(signum, act, oldact);
get_chained_signals().at(signum) = chained_siginfo{signum, nullptr, *act};
// make sure rocprofv3_error_signal_handler doesn't call itself
if((act->sa_flags & SA_SIGINFO) == SA_SIGINFO &&
act->sa_sigaction != &rocprofv3_error_signal_handler)
get_chained_signals().at(signum) = chained_siginfo{signum, nullptr, *act};
struct sigaction _upd_act = *act;
_upd_act.sa_flags |= (SA_SIGINFO | SA_RESETHAND | SA_NOCLDSTOP);
@@ -2959,6 +3008,7 @@ rocprofv3_main(int argc, char** argv, char** envp)
auto _data = std::vector<std::string_view>{};
size_t n = 0;
const char* p = nullptr;
if(!inp) return _data;
do
{
p = inp[n++];
@@ -2980,10 +3030,26 @@ rocprofv3_main(int argc, char** argv, char** envp)
ROCP_INFO << "rocprofv3: main function wrapper will be invoked...";
auto _main_timer = std::optional<common::simple_timer>{};
// should never happen but if it does, don't time
if(!_argv.empty())
_main_timer = common::simple_timer{
fmt::format("[rocprofv3] '{}'", fmt::join(_argv.begin(), _argv.end(), " "))};
if(tool_metadata && tool_metadata->process_start_ns == 0)
rocprofiler_get_timestamp(&(tool_metadata->process_start_ns));
auto ret = CHECK_NOTNULL(get_main_function())(argc, argv, envp);
if(tool_metadata && tool_metadata->process_end_ns == 0)
rocprofiler_get_timestamp(&(tool_metadata->process_end_ns));
ROCP_INFO << "rocprofv3: main function has returned with exit code: " << ret;
// reset so that it reports the timing
if(_main_timer) _main_timer.reset();
finalize_rocprofv3(__FUNCTION__);
ROCP_INFO << "rocprofv3 finished. exit code: " << ret;
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp
+3 -3
Просмотреть файл
@@ -294,9 +294,9 @@ TEST(dimension, block_dim_test)
rocprofiler_iterate_counter_dimensions(
{.handle = metric.id()},
[](rocprofiler_counter_id_t,
const rocprofiler_record_dimension_info_t* dim_info,
size_t num_dims,
void* user_data) -> rocprofiler_status_t {
const rocprofiler_counter_record_dimension_info_t* dim_info,
size_t num_dims,
void* user_data) -> rocprofiler_status_t {
auto expected_dims = *static_cast<
std::unordered_map<counters::rocprofiler_profile_counter_instance_types,
uint64_t>*>(user_data);
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/stream.cpp
+32 -18
Просмотреть файл
@@ -185,9 +185,11 @@ FuncT create_write_functor(RetT (*func)(Args...))
external_corr_ids);
assert(buffered_contexts.empty() && "Stream tracing should not have any buffered contexts");
auto tracer_data = common::init_public_api_struct(callback_api_data_t{});
auto internal_corr_id = 0;
auto ancestor_corr_id = 0;
auto tracer_data = common::init_public_api_struct(callback_api_data_t{},
rocprofiler_stream_id_t{.handle = 0},
rocprofiler_address_t{.value = 0});
constexpr auto stream_idx = common::mpl::index_of<hipStream_t*, function_args_type>::value;
auto stream = std::get<stream_idx>(std::make_tuple(std::forward<Args>(args)...));
@@ -200,7 +202,8 @@ FuncT create_write_functor(RetT (*func)(Args...))
{
if(stream)
{
tracer_data.stream_id = add_stream(*stream);
tracer_data.stream_id = add_stream(*stream);
tracer_data.stream_value.ptr = *stream;
}
tracing::execute_phase_none_callbacks(callback_contexts,
thr_id,
@@ -248,9 +251,11 @@ FuncT create_destroy_functor(RetT (*func)(Args...))
external_corr_ids);
assert(buffered_contexts.empty() && "Stream tracing should not have any buffered contexts");
auto tracer_data = common::init_public_api_struct(callback_api_data_t{});
auto internal_corr_id = 0;
auto ancestor_corr_id = 0;
auto tracer_data = common::init_public_api_struct(callback_api_data_t{},
rocprofiler_stream_id_t{.handle = 0},
rocprofiler_address_t{.value = 0});
auto stream = std::get<stream_idx>(std::make_tuple(std::forward<Args>(args)...));
@@ -261,7 +266,8 @@ FuncT create_destroy_functor(RetT (*func)(Args...))
if(!callback_contexts.empty())
{
tracer_data.stream_id = get_stream_id(stream);
tracer_data.stream_id = get_stream_id(stream);
tracer_data.stream_value.ptr = stream;
tracing::execute_phase_none_callbacks(callback_contexts,
thr_id,
internal_corr_id,
@@ -309,15 +315,18 @@ FuncT create_read_functor(RetT (*func)(Args...))
assert(buffered_contexts.empty() && "Stream tracing should not have any buffered contexts");
auto tracer_data = common::init_public_api_struct(callback_api_data_t{});
auto internal_corr_id = 0;
auto ancestor_corr_id = 0;
auto tracer_data = common::init_public_api_struct(callback_api_data_t{},
rocprofiler_stream_id_t{.handle = 0},
rocprofiler_address_t{.value = 0});
auto stream = std::get<stream_idx>(std::make_tuple(std::forward<Args>(args)...));
if(!callback_contexts.empty())
{
tracer_data.stream_id = get_stream_id(stream);
tracer_data.stream_id = get_stream_id(stream);
tracer_data.stream_value.ptr = stream;
tracing::execute_phase_enter_callbacks(callback_contexts,
thr_id,
internal_corr_id,
@@ -385,6 +394,9 @@ update_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
using info_type = hip_api_info<TableIdx, OpIdx>;
using function_args_type = decltype(info_type::get_args_type());
static_assert(info_type::table_idx == ROCPROFILER_HIP_TABLE_ID_Runtime,
"This function should only be instantiated for HIP runtime API");
if constexpr(std::is_same<table_type, Tp>::value &&
(common::mpl::is_one_of<hipStream_t, function_args_type>::value ||
common::mpl::is_one_of<hipStream_t*, function_args_type>::value))
@@ -394,7 +406,7 @@ update_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
// make sure we don't access a field that doesn't exist in input table
if(_info.offset() >= _orig->size) return;
ROCP_TRACE << "updating table entry for " << _info.name;
ROCP_TRACE << fmt::format("[hip stream] checking table entry for {}", _info.name);
constexpr auto num_args = function_args_type::size();
@@ -404,6 +416,9 @@ update_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
common::mpl::index_of<hipStream_t, function_args_type>::value;
constexpr auto rstream_idx =
common::mpl::index_of<hipStream_t, common::mpl::reverse<function_args_type>>::value;
constexpr auto is_hip_stream_destroy_func =
info_type::table_idx == ROCPROFILER_HIP_TABLE_ID_Runtime &&
info_type::operation_idx == ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamDestroy;
// index_of finds the first argument of that type. So find the first and last
// arg of the given type and make sure it resolves to the same distance
@@ -413,21 +428,19 @@ update_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
// 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);
constexpr auto is_hip_destroy_func =
std::is_same<decltype(info_type::operation_idx),
rocprofiler_hip_runtime_api_id_t>::value &&
(static_cast<rocprofiler_hip_runtime_api_id_t>(info_type::operation_idx) ==
ROCPROFILER_HIP_RUNTIME_API_ID_hipStreamDestroy);
if constexpr(is_hip_destroy_func)
auto& _table = _info.get_table(_orig);
auto& _func = _info.get_table_func(_table);
if constexpr(is_hip_stream_destroy_func)
{
ROCP_INFO << _info.name << " has been designated as a stream destroy function";
ROCP_INFO << fmt::format(
"[hip stream] {} has been designated as a stream destroy function", _info.name);
_func = create_destroy_functor<TableIdx, OpIdx>(_func);
}
else
{
ROCP_INFO << _info.name << " has been designated as a stream set function";
ROCP_INFO << fmt::format(
"[hip stream] {} has been designated as a stream set function", _info.name);
_func = create_read_functor<TableIdx, OpIdx>(_func);
}
}
@@ -444,7 +457,8 @@ update_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
static_assert(stream_idx == (num_args - rstream_idx - 1),
"function has more than one stream argument");
ROCP_INFO << _info.name << " has been designated as a stream create function";
ROCP_INFO << fmt::format(
"[hip stream] {} has been designated as a stream create function", _info.name);
// 1. get the sub-table containing the function pointer in original table
// 2. get reference to function pointer in sub-table in original table
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/stream.hpp
-3
Просмотреть файл
@@ -38,9 +38,6 @@ namespace stream
{
using hip_runtime_api_table_t = HipDispatchTable;
rocprofiler_stream_id_t
get_stream_id();
const char*
name_by_id(uint32_t id);
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/intercept_table.cpp
+56
Просмотреть файл
@@ -50,6 +50,29 @@ namespace
template <rocprofiler_intercept_table_t... Idx>
using library_sequence_t = std::integer_sequence<rocprofiler_intercept_table_t, Idx...>;
#define ROCPROFILER_INTERCEPT_TABLE_KIND_STRING(TABLE, NAME) \
template <> \
struct intercept_table_info<ROCPROFILER_##TABLE##_TABLE> \
{ \
static constexpr auto value = ROCPROFILER_##TABLE##_TABLE; \
static constexpr auto name = NAME; \
static constexpr auto length = std::string_view{NAME}.length(); \
static constexpr auto data = std::pair<const char*, size_t>{name, length}; \
};
template <rocprofiler_intercept_table_t Idx>
struct intercept_table_info;
ROCPROFILER_INTERCEPT_TABLE_KIND_STRING(HSA, "HSA")
ROCPROFILER_INTERCEPT_TABLE_KIND_STRING(HIP_RUNTIME, "HIP (runtime)")
ROCPROFILER_INTERCEPT_TABLE_KIND_STRING(HIP_COMPILER, "HIP (compiler)")
ROCPROFILER_INTERCEPT_TABLE_KIND_STRING(MARKER_CORE, "MARKER (ROCTx)")
ROCPROFILER_INTERCEPT_TABLE_KIND_STRING(MARKER_CONTROL, "MARKER (ROCTx Control)")
ROCPROFILER_INTERCEPT_TABLE_KIND_STRING(MARKER_NAME, "MARKER (ROCTx Name)")
ROCPROFILER_INTERCEPT_TABLE_KIND_STRING(RCCL, "RCCL")
ROCPROFILER_INTERCEPT_TABLE_KIND_STRING(ROCDECODE, "rocDecode")
ROCPROFILER_INTERCEPT_TABLE_KIND_STRING(ROCJPEG, "rocJPEG")
// this is used to loop over the different libraries
constexpr auto intercept_library_seq = library_sequence_t<ROCPROFILER_HSA_TABLE,
ROCPROFILER_HIP_RUNTIME_TABLE,
@@ -85,6 +108,25 @@ get_intercept()
return _v;
}
// adds callbacks to intercept instance(s)
template <rocprofiler_intercept_table_t Idx, rocprofiler_intercept_table_t... Tail>
auto
get_intercept_table_info(rocprofiler_intercept_table_t kind, library_sequence_t<Idx, Tail...>)
{
using info_type = intercept_table_info<Idx>;
using return_type = decltype(info_type::data);
if(info_type::value == kind)
{
return info_type::data;
}
if constexpr(sizeof...(Tail) > 0)
return get_intercept_table_info(kind, library_sequence_t<Tail...>{});
return return_type{nullptr, 0};
}
// adds callbacks to intercept instance(s)
template <rocprofiler_intercept_table_t... Idx>
void
@@ -208,6 +250,20 @@ template void notify_intercept_table_registration(rocprofiler_intercept_table_t,
} // namespace rocprofiler
extern "C" {
rocprofiler_status_t
rocprofiler_query_intercept_table_name(rocprofiler_intercept_table_t kind,
const char** name,
uint64_t* name_len)
{
auto&& val = rocprofiler::intercept_table::get_intercept_table_info(
kind, rocprofiler::intercept_table::intercept_library_seq);
if(name) *name = val.first;
if(name_len) *name_len = val.second;
return (val.first) ? ROCPROFILER_STATUS_SUCCESS : ROCPROFILER_STATUS_ERROR_KIND_NOT_FOUND;
}
rocprofiler_status_t
rocprofiler_at_intercept_table_registration(rocprofiler_intercept_library_cb_t callback,
int libs,
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/registration.cpp
+5 -3
Просмотреть файл
@@ -28,6 +28,7 @@
#include "lib/common/filesystem.hpp"
#include "lib/common/logging.hpp"
#include "lib/common/static_object.hpp"
#include "lib/common/static_tl_object.hpp"
#include "lib/rocprofiler-sdk/agent.hpp"
#include "lib/rocprofiler-sdk/code_object/code_object.hpp"
#include "lib/rocprofiler-sdk/context/context.hpp"
@@ -244,12 +245,12 @@ find_clients()
return true;
};
constexpr auto client_id_size = sizeof(rocprofiler_client_id_t);
auto emplace_client = [&data, priority_offset](
auto emplace_client = [&data, priority_offset](
std::string_view _name,
void* _dlhandle,
auto* _cfg_func) -> std::optional<client_library>& {
uint32_t _prio = priority_offset + data.size();
constexpr auto client_id_size = sizeof(rocprofiler_client_id_t);
uint32_t _prio = priority_offset + data.size();
return data.emplace_back(
client_library{std::string{_name},
_dlhandle,
@@ -680,6 +681,7 @@ initialize()
set_init_status(-1);
std::atexit([]() {
finalize();
common::destroy_static_tl_objects();
common::destroy_static_objects();
});
init_logging();
projects/rocprofiler-sdk/source/lib/tests/common/CMakeLists.txt
+1 -1
Просмотреть файл
@@ -5,7 +5,7 @@ project(rocprofiler-sdk-tests-common LANGUAGES C CXX)
include(GoogleTest)
set(common_sources demangling.cpp environment.cpp mpl.cpp)
set(common_sources demangling.cpp environment.cpp md5sum.cpp mpl.cpp)
add_executable(common-tests)
target_sources(common-tests PRIVATE ${common_sources})
projects/rocprofiler-sdk/source/lib/tests/common/md5sum.cpp
+58
Просмотреть файл
@@ -0,0 +1,58 @@
// MIT License
//
// Copyright (c) 2023-2025 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/common/md5sum.hpp"
#include <fmt/format.h>
#include <gtest/gtest.h>
#include <cstddef>
#include <cstdint>
#include <sstream>
#include <string>
TEST(common, md5sum)
{
auto _val = rocprofiler::common::md5sum{};
_val.update(234234);
_val.finalize();
auto _hex_digest = _val.hexdigest();
EXPECT_EQ(_hex_digest, "5cdf897625070d805f3fb2469b08eebb");
auto read_hex = [](std::string&& _inp) {
int _ret = 0;
std::stringstream{_inp} >> std::hex >> _ret;
return _ret;
};
auto _raw_digest = _val.rawdigest();
for(size_t i = 0; i < _raw_digest.size(); ++i)
{
auto _sub = _hex_digest.substr(i * 2, 2);
auto _extract = read_hex(fmt::format("0x{}", _sub));
int _raw = _raw_digest.at(i);
EXPECT_EQ(_raw, _extract) << fmt::format(
"i={}, raw[i]={}, hex={}, hexdigest={}", i, _raw, _sub, _hex_digest);
}
}
projects/rocprofiler-sdk/source/scripts/install-apt-deps.sh
+1 -1
Просмотреть файл
@@ -1,7 +1,7 @@
#!/bin/bash -e
apt-get update
apt-get install -y cmake clang-tidy-15 g++-11 g++-12 python3-pip libdw-dev
apt-get install -y cmake clang-tidy-15 g++-11 g++-12 python3-pip libdw-dev libsqlite3-dev
update-alternatives --install /usr/bin/clang-tidy clang-tidy /usr/bin/clang-tidy-17 10
update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-11 10 --slave /usr/bin/g++ g++ /usr/bin/g++-11
update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-12 20 --slave /usr/bin/g++ g++ /usr/bin/g++-12
projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py
+14 -5
Просмотреть файл
@@ -26,6 +26,15 @@ import sys
import pytest
test_api_traces = [
"hsa_api_traces",
"marker_api_traces",
"hip_api_traces",
"rccl_api_traces",
"scratch_memory_traces",
]
# helper function
def node_exists(name, data, min_len=1):
assert name in data
@@ -153,7 +162,7 @@ def test_timestamps(input_data):
cb_start = {}
cb_end = {}
for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]:
for titr in test_api_traces:
for itr in sdk_data["callback_records"][titr]:
cid = itr["correlation_id"]["internal"]
phase = itr["phase"]
@@ -197,7 +206,7 @@ def test_internal_correlation_ids(input_data):
sdk_data = data["rocprofiler-sdk-json-tool"]
api_corr_ids = []
for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]:
for titr in test_api_traces:
for itr in sdk_data["callback_records"][titr]:
api_corr_ids.append(itr["correlation_id"]["internal"])
@@ -223,14 +232,14 @@ def test_external_correlation_ids(input_data):
sdk_data = data["rocprofiler-sdk-json-tool"]
extern_corr_ids = []
for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]:
for titr in test_api_traces:
for itr in sdk_data["callback_records"][titr]:
assert itr["correlation_id"]["external"] > 0
assert itr["thread_id"] == itr["correlation_id"]["external"]
extern_corr_ids.append(itr["correlation_id"]["external"])
extern_corr_ids = list(set(sorted(extern_corr_ids)))
for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]:
for titr in test_api_traces:
for itr in sdk_data["buffer_records"][titr]:
assert itr["correlation_id"]["external"] > 0, f"[{titr}] {itr}"
assert (
@@ -433,7 +442,7 @@ def test_retired_correlation_ids(input_data):
return dict(sorted(inp.items()))
api_corr_ids = {}
for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]:
for titr in test_api_traces:
for itr in sdk_data["buffer_records"][titr]:
corr_id = itr["correlation_id"]["internal"]
name = get_operation_name(buffer_records, itr["kind"], itr["operation"])
projects/rocprofiler-sdk/tests/bin/hip-graph/hip-graph.cpp
+8 -2
Просмотреть файл
@@ -153,8 +153,8 @@ run(uint64_t devid,
for(uint64_t i = 0; i < nstream; ++i)
checkHipErrors(hipGraphLaunch(execs.at(i), streams.at(i)));
log_message("synchronizing device");
checkHipErrors(hipDeviceSynchronize());
for(uint64_t i = 0; i < nstream; ++i)
checkHipErrors(hipStreamSynchronize(streams.at(i)));
log_message("destroying graph");
for(uint64_t i = 0; i < nstream; ++i)
@@ -207,6 +207,12 @@ main(int argc, char* argv[])
for(auto& itr : threads)
itr.join();
for(uint64_t i = 0; i < ndevice; ++i)
{
checkHipErrors(hipSetDevice(i));
checkHipErrors(hipDeviceSynchronize());
}
std::cout << "[" << ::basename(argv[0]) << "] complete" << std::endl;
return 0;
}
projects/rocprofiler-sdk/tests/bin/page-migration/page-migration.cpp
+17 -5
Просмотреть файл
@@ -118,7 +118,7 @@ private:
};
int
run_test(int num_iter)
run_test(int dev_idx, int num_iter)
{
using namespace std::chrono_literals;
@@ -141,6 +141,7 @@ run_test(int num_iter)
elem_count,
data_v);
HIP_API_CALL(hipSetDevice(dev_idx));
HIP_API_CALL(hipHostRegister(data, elem_count * sizeof(size_t), hipHostRegisterDefault));
constexpr size_t MAPS_BUFFER_SIZE = 1024 * 1024;
@@ -221,8 +222,9 @@ run_test(int num_iter)
}
}
HIP_API_CALL(hipStreamSynchronize(stream));
HIP_API_CALL(hipStreamDestroy(stream));
HIP_API_CALL(hipDeviceSynchronize());
HIP_API_CALL(hipHostUnregister(data));
return 0;
}
@@ -266,16 +268,19 @@ main(int argc, const char** argv)
exit(EXIT_FAILURE);
}
run_test(num_iter);
int ndevice = 0;
HIP_API_CALL(hipGetDeviceCount(&ndevice));
run_test(0, num_iter);
std::vector<std::thread> threads;
threads.reserve(num_threads);
std::cerr << "Running " << num_iter << " iterations/thread on " << num_threads << " threads\n";
for(auto i = 0; i < num_threads; ++i)
{
threads.emplace_back([_num_iter = num_iter]() { run_test(_num_iter); });
auto _dev_idx = i % ndevice;
threads.emplace_back([_dev_idx, _num_iter = num_iter]() { run_test(_dev_idx, _num_iter); });
}
std::cerr << "Waiting for threads\n";
@@ -284,5 +289,12 @@ main(int argc, const char** argv)
t.join();
}
for(int i = 0; i < ndevice; ++i)
{
HIP_API_CALL(hipSetDevice(i));
HIP_API_CALL(hipDeviceSynchronize());
HIP_API_CALL(hipDeviceReset());
}
return 0;
}
projects/rocprofiler-sdk/tests/kernel-tracing/validate.py
+13 -5
Просмотреть файл
@@ -25,6 +25,14 @@
import sys
import pytest
test_api_traces = [
"hsa_api_traces",
"marker_api_traces",
"hip_api_traces",
"rccl_api_traces",
"scratch_memory_traces",
]
# helper function
def node_exists(name, data, min_len=1):
@@ -103,7 +111,7 @@ def test_timestamps(input_data):
cb_start = {}
cb_end = {}
for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]:
for titr in test_api_traces:
for itr in sdk_data["callback_records"][titr]:
cid = itr["correlation_id"]["internal"]
phase = itr["phase"]
@@ -161,7 +169,7 @@ def test_internal_correlation_ids(input_data):
sdk_data = data["rocprofiler-sdk-json-tool"]
api_corr_ids = []
for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]:
for titr in test_api_traces:
for itr in sdk_data["callback_records"][titr]:
api_corr_ids.append(itr["correlation_id"]["internal"])
@@ -190,7 +198,7 @@ def test_retired_correlation_ids(input_data):
return dict(sorted(inp.items()))
api_corr_ids = {}
for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]:
for titr in test_api_traces:
for itr in sdk_data["buffer_records"][titr]:
corr_id = itr["correlation_id"]["internal"]
assert corr_id not in api_corr_ids.keys()
@@ -233,14 +241,14 @@ def test_external_correlation_ids(input_data):
sdk_data = data["rocprofiler-sdk-json-tool"]
extern_corr_ids = []
for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]:
for titr in test_api_traces:
for itr in sdk_data["callback_records"][titr]:
assert itr["correlation_id"]["external"] > 0
assert itr["thread_id"] == itr["correlation_id"]["external"]
extern_corr_ids.append(itr["correlation_id"]["external"])
extern_corr_ids = list(set(sorted(extern_corr_ids)))
for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]:
for titr in test_api_traces:
for itr in sdk_data["buffer_records"][titr]:
assert itr["correlation_id"]["external"] > 0, f"[{titr}] {itr}"
assert (
projects/rocprofiler-sdk/tests/page-migration/CMakeLists.txt
+1
Просмотреть файл
@@ -30,6 +30,7 @@ set(page-migration-env
"${PRELOAD_ENV}"
"ROCPROFILER_DISABLE_PERFETTO_ANNOTATIONS=1"
"ROCPROFILER_TOOL_OUTPUT_FILE=page-migration-test.json"
"ROCPROFILER_TOOL_CONTEXTS_EXCLUDE=HSA_API_BUFFERED,HSA_API_CALLBACK"
"LD_LIBRARY_PATH=$<TARGET_FILE_DIR:rocprofiler-sdk::rocprofiler-sdk-shared-library>:$ENV{LD_LIBRARY_PATH}"
)
projects/rocprofiler-sdk/tests/page-migration/validate.py
+17 -9
Просмотреть файл
@@ -23,11 +23,19 @@
# THE SOFTWARE.
from collections import defaultdict
import os
import sys
import pytest
test_api_traces = [
"hsa_api_traces",
"marker_api_traces",
"hip_api_traces",
"rccl_api_traces",
"scratch_memory_traces",
]
# helper function
def node_exists(name, data, min_len=1):
assert name in data
@@ -105,15 +113,15 @@ def test_data_structure(input_data):
node_exists("code_objects", sdk_data["callback_records"])
node_exists("kernel_symbols", sdk_data["callback_records"])
node_exists("host_functions", sdk_data["callback_records"])
node_exists("hsa_api_traces", sdk_data["callback_records"])
node_exists("hip_api_traces", sdk_data["callback_records"], 0)
node_exists("hsa_api_traces", sdk_data["callback_records"], 0)
node_exists("hip_api_traces", sdk_data["callback_records"])
node_exists("marker_api_traces", sdk_data["callback_records"], 0)
node_exists("names", sdk_data["buffer_records"])
node_exists("kernel_dispatch", sdk_data["buffer_records"])
node_exists("memory_copies", sdk_data["buffer_records"], 0)
node_exists("hsa_api_traces", sdk_data["buffer_records"])
node_exists("hip_api_traces", sdk_data["buffer_records"], 0)
node_exists("hsa_api_traces", sdk_data["buffer_records"], 0)
node_exists("hip_api_traces", sdk_data["buffer_records"])
node_exists("marker_api_traces", sdk_data["buffer_records"], 0)
node_exists("retired_correlation_ids", sdk_data["buffer_records"])
node_exists("page_migration", sdk_data["buffer_records"])
@@ -125,7 +133,7 @@ def test_timestamps(input_data):
cb_start = {}
cb_end = {}
for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]:
for titr in test_api_traces:
for itr in sdk_data["callback_records"][titr]:
cid = itr["correlation_id"]["internal"]
phase = itr["phase"]
@@ -161,7 +169,7 @@ def test_internal_correlation_ids(input_data):
sdk_data = data["rocprofiler-sdk-json-tool"]
api_corr_ids = []
for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]:
for titr in test_api_traces:
for itr in sdk_data["callback_records"][titr]:
api_corr_ids.append(itr["correlation_id"]["internal"])
@@ -187,14 +195,14 @@ def test_external_correlation_ids(input_data):
sdk_data = data["rocprofiler-sdk-json-tool"]
extern_corr_ids = []
for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]:
for titr in test_api_traces:
for itr in sdk_data["callback_records"][titr]:
assert itr["correlation_id"]["external"] > 0
assert itr["thread_id"] == itr["correlation_id"]["external"]
extern_corr_ids.append(itr["correlation_id"]["external"])
extern_corr_ids = list(set(sorted(extern_corr_ids)))
for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]:
for titr in test_api_traces:
for itr in sdk_data["buffer_records"][titr]:
assert itr["correlation_id"]["external"] > 0
assert itr["thread_id"] == itr["correlation_id"]["external"]
projects/rocprofiler-sdk/tests/pytest-packages/tests/rocprofv3.py
+14 -1
Просмотреть файл
@@ -66,14 +66,27 @@ def test_perfetto_data(
else:
unique_counter_ids = set()
counter_info = {}
for itr in json_data["rocprofiler-sdk-tool"]["counters"]:
counter_info[itr["id"]["handle"]] = itr
agent_info = {}
for itr in json_data["rocprofiler-sdk-tool"]["agents"]:
agent_info[itr["id"]["handle"]] = itr
for dispatch_entry in json_data["rocprofiler-sdk-tool"]["callback_records"][
js_category
]:
counter_records = dispatch_entry["records"]
agent_id = dispatch_entry["dispatch_data"]["dispatch_info"]["agent_id"][
"handle"
]
agent_node_id = agent_info[agent_id]["node_id"]
for record in counter_records:
counter_id = record["counter_id"]["handle"]
unique_counter_ids.add(counter_id)
counter_name = counter_info[counter_id]["name"]
unique_counter_ids.add(f"Agent [{agent_node_id}] PMC {counter_name}")
_js_data = [{"counter_id": id} for id in unique_counter_ids]
projects/rocprofiler-sdk/tests/rocprofv3/kernel-rename/validate.py
+1 -1
Просмотреть файл
@@ -140,7 +140,7 @@ def test_kernel_trace(json_data):
external_corr_id = dispatch["correlation_id"]["external"]
assert external_corr_id > 0
kernel_rename = get_kernel_rename(dispatch.kernel_rename_val)
kernel_rename = get_kernel_rename(external_corr_id)
assert kernel_rename is not None, f"{dispatch}"
assert kernel_rename != kernel_name, f"{dispatch}"
assert (
projects/rocprofiler-sdk/tests/rocprofv3/summary/validate.py
+1 -1
Просмотреть файл
@@ -117,7 +117,7 @@ def test_kernel_trace(json_data):
external_corr_id = dispatch["correlation_id"]["external"]
assert external_corr_id > 0
kernel_rename = get_kernel_rename(dispatch.kernel_rename_val)
kernel_rename = get_kernel_rename(external_corr_id)
assert kernel_rename is not None, f"{dispatch}"
assert kernel_rename != kernel_name, f"{dispatch}"
assert (
projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-counter-collection/CMakeLists.txt
+10 -10
Просмотреть файл
@@ -14,13 +14,13 @@ find_package(rocprofiler-sdk REQUIRED)
rocprofiler_configure_pytest_files(COPY validate.py conftest.py input.txt
CONFIG pytest.ini)
# pmc3
# pmc
add_test(
NAME rocprofv3-test-tracing-plus-counter-collection-execute
COMMAND
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --hsa-trace -i
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --hsa-trace --kernel-trace -i
${CMAKE_CURRENT_BINARY_DIR}/input.txt -d
${CMAKE_CURRENT_BINARY_DIR}/out_counter_collection_trace -o pmc3 --output-format
${CMAKE_CURRENT_BINARY_DIR}/out_counter_collection_trace -o pmc --output-format
JSON PFTRACE CSV -- $<TARGET_FILE:simple-transpose>)
string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV
@@ -39,15 +39,15 @@ foreach(_DIR "pmc_1" "pmc_2" "pmc_3" "pmc_4")
NAME rocprofv3-test-tracing-plus-counter-collection-validate-${_DIR}
COMMAND
${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --json-input
"${CMAKE_CURRENT_BINARY_DIR}/out_counter_collection_trace/${_DIR}/pmc3_results.json"
"${CMAKE_CURRENT_BINARY_DIR}/out_counter_collection_trace/${_DIR}/pmc_results.json"
--pftrace-input
"${CMAKE_CURRENT_BINARY_DIR}/out_counter_collection_trace/${_DIR}/pmc3_results.pftrace"
"${CMAKE_CURRENT_BINARY_DIR}/out_counter_collection_trace/${_DIR}/pmc_results.pftrace"
--hsa-input
"${CMAKE_CURRENT_BINARY_DIR}/out_counter_collection_trace/${_DIR}/pmc3_hsa_api_trace.csv"
"${CMAKE_CURRENT_BINARY_DIR}/out_counter_collection_trace/${_DIR}/pmc_hsa_api_trace.csv"
--agent-input
"${CMAKE_CURRENT_BINARY_DIR}/out_counter_collection_trace/${_DIR}/pmc3_agent_info.csv"
"${CMAKE_CURRENT_BINARY_DIR}/out_counter_collection_trace/${_DIR}/pmc_agent_info.csv"
--counter-input
"${CMAKE_CURRENT_BINARY_DIR}/out_counter_collection_trace/${_DIR}/pmc3_counter_collection.csv"
"${CMAKE_CURRENT_BINARY_DIR}/out_counter_collection_trace/${_DIR}/pmc_counter_collection.csv"
)
set_tests_properties(
@@ -68,10 +68,10 @@ foreach(_GROUP "single" "multiple")
add_test(
NAME ${_TEST_NAME_PREFIX}-execute
COMMAND
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --hip-trace --hsa-trace --pmc
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --sys-trace --kernel-rename --pmc
${cmdl_${_GROUP}} -d ${CMAKE_CURRENT_BINARY_DIR}/out_counter_collection_trace
-o cmdl/${_GROUP} --output-format JSON PFTRACE CSV --
$<TARGET_FILE:simple-transpose>)
$<TARGET_FILE:transpose> 2 50 5)
set_tests_properties(
${_TEST_NAME_PREFIX}-execute
projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-counter-collection/validate.py
+61
Просмотреть файл
@@ -22,6 +22,7 @@
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
# THE SOFTWARE.
import os
import sys
import pytest
@@ -44,6 +45,66 @@ def test_validate_counter_collection_plus_tracing(
assert len(counter_collection_data) > 0
def test_validate_counter_collection_plus_tracing_dispatch_data(json_data):
"""
Validates that the dispatch data from kernel tracing and counter collection match
"""
data = json_data["rocprofiler-sdk-tool"]
argv0 = os.path.basename(data["metadata"]["command"][0])
kernel_ids = []
kernel_dispatch_data = {}
counter_collection_data = {}
for itr in data["buffer_records"]["kernel_dispatch"]:
kern_id = itr["dispatch_info"]["kernel_id"]
kernel_dispatch_data[kern_id] = itr
kernel_dispatch_data[kern_id]["stream_id"] = itr["stream_id"]["handle"]
if kern_id not in kernel_ids:
kernel_ids += [kern_id]
for itr in data["callback_records"]["counter_collection"]:
kern_id = itr["dispatch_data"]["dispatch_info"]["kernel_id"]
counter_collection_data[kern_id] = itr["dispatch_data"]
counter_collection_data[kern_id]["thread_id"] = itr["thread_id"]
counter_collection_data[kern_id]["stream_id"] = itr["stream_id"]["handle"]
if kern_id not in kernel_ids:
kernel_ids += [kern_id]
kernel_ids = sorted(kernel_ids)
for kern_id in kernel_ids:
for itr in ["thread_id", "stream_id", "start_timestamp", "end_timestamp"]:
assert (
counter_collection_data[kern_id][itr]
== kernel_dispatch_data[kern_id][itr]
), f"kernel_id={kern_id}, check={itr}"
if itr == "stream_id":
if argv0 == "simple-transpose":
assert (
kernel_dispatch_data[kern_id][itr] == 0
), f"kernel_id={kern_id}, check={itr}"
elif argv0 == "transpose":
assert (
kernel_dispatch_data[kern_id][itr] > 0
), f"kernel_id={kern_id}, check={itr}"
else:
assert False, f"no stream_id handling for '{argv0}'"
else:
assert (
kernel_dispatch_data[kern_id][itr] > 0
), f"kernel_id={kern_id}, check={itr}"
assert (
counter_collection_data[kern_id]["correlation_id"]["internal"]
== kernel_dispatch_data[kern_id]["correlation_id"]["internal"]
)
assert (
counter_collection_data[kern_id]["correlation_id"]["external"]
== kernel_dispatch_data[kern_id]["correlation_id"]["external"]
)
def test_perfetto_data(pftrace_data, json_data):
import rocprofiler_sdk.tests.rocprofv3 as rocprofv3
projects/rocprofiler-sdk/tests/tools/json-tool.cpp
+19 -1
Просмотреть файл
@@ -966,7 +966,8 @@ tool_tracing_buffered(rocprofiler_context_id_t /*context*/,
marker_api_bf_records.emplace_back(*record);
}
else if(header->kind == ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API)
else if(header->kind == ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API ||
header->kind == ROCPROFILER_BUFFER_TRACING_HIP_COMPILER_API)
{
auto* record =
static_cast<rocprofiler_buffer_tracing_hip_api_record_t*>(header->payload);
@@ -1309,6 +1310,15 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
nullptr),
"hip runtime api callback tracing service configure");
// ROCPROFILER_CALL(rocprofiler_configure_callback_tracing_service(
// hip_api_callback_ctx,
// ROCPROFILER_CALLBACK_TRACING_HIP_COMPILER_API,
// nullptr,
// 0,
// tool_tracing_callback,
// nullptr),
// "hip compiler api callback tracing service configure");
ROCPROFILER_CALL(
rocprofiler_configure_callback_tracing_service(code_object_ctx,
ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT,
@@ -1592,6 +1602,14 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
hip_api_buffered_buffer),
"buffer tracing service configure");
// ROCPROFILER_CALL(
// rocprofiler_configure_buffer_tracing_service(hip_api_buffered_ctx,
// ROCPROFILER_BUFFER_TRACING_HIP_COMPILER_API,
// nullptr,
// 0,
// hip_api_buffered_buffer),
// "buffer tracing service configure");
ROCPROFILER_CALL(
rocprofiler_configure_buffer_tracing_service(marker_api_buffered_ctx,
ROCPROFILER_BUFFER_TRACING_MARKER_CORE_API,