diff --git a/.github/workflows/codeql.yml b/.github/workflows/codeql.yml index deaefbb332..39298c2869 100644 --- a/.github/workflows/codeql.yml +++ b/.github/workflows/codeql.yml @@ -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 diff --git a/.github/workflows/continuous_integration.yml b/.github/workflows/continuous_integration.yml index a788feba4d..9ab5824359 100644 --- a/.github/workflows/continuous_integration.yml +++ b/.github/workflows/continuous_integration.yml @@ -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* diff --git a/.github/workflows/docs.yml b/.github/workflows/docs.yml index 4decdfc94b..3a651c47e9 100644 --- a/.github/workflows/docs.yml +++ b/.github/workflows/docs.yml @@ -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 diff --git a/.github/workflows/generate_afar.yml b/.github/workflows/generate_afar.yml index b9dbfe28dd..e789ae7683 100644 --- a/.github/workflows/generate_afar.yml +++ b/.github/workflows/generate_afar.yml @@ -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 diff --git a/.gitmodules b/.gitmodules index fad932f436..7bbab5a350 100644 --- a/.gitmodules +++ b/.gitmodules @@ -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 diff --git a/CMakeLists.txt b/CMakeLists.txt index 53f9e247ad..bc1daaf1eb 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -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 diff --git a/cmake/rocprofiler_build_settings.cmake b/cmake/rocprofiler_build_settings.cmake index 3abdbed653..24c34bc0cd 100644 --- a/cmake/rocprofiler_build_settings.cmake +++ b/cmake/rocprofiler_build_settings.cmake @@ -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 # diff --git a/cmake/rocprofiler_interfaces.cmake b/cmake/rocprofiler_interfaces.cmake index 45b15042c8..df604b80ff 100644 --- a/cmake/rocprofiler_interfaces.cmake +++ b/cmake/rocprofiler_interfaces.cmake @@ -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) diff --git a/cmake/rocprofiler_options.cmake b/cmake/rocprofiler_options.cmake index 9b48a0472b..dfcf140672 100644 --- a/cmake/rocprofiler_options.cmake +++ b/cmake/rocprofiler_options.cmake @@ -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) diff --git a/docker/Dockerfile b/docker/Dockerfile index 2ac578e224..66ba89d411 100644 --- a/docker/Dockerfile +++ b/docker/Dockerfile @@ -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 && \ diff --git a/external/CMakeLists.txt b/external/CMakeLists.txt index 4298bed65a..e864ded3f2 100644 --- a/external/CMakeLists.txt +++ b/external/CMakeLists.txt @@ -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 $) target_include_directories( rocprofiler-sdk-json SYSTEM INTERFACE $) + +# +# 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 + /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 + $ + ) + target_include_directories( + rocprofiler-sdk-sqlite3 SYSTEM + INTERFACE $ + ) + 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 $) +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) diff --git a/external/gotcha b/external/gotcha new file mode 160000 index 0000000000..9afbaf0ae5 --- /dev/null +++ b/external/gotcha @@ -0,0 +1 @@ +Subproject commit 9afbaf0ae5858098d50f79bf494b6333da17da6b diff --git a/external/pybind11 b/external/pybind11 new file mode 160000 index 0000000000..914c06fb25 --- /dev/null +++ b/external/pybind11 @@ -0,0 +1 @@ +Subproject commit 914c06fb252b6cc3727d0eedab6736e88a3fcb01 diff --git a/external/sqlite b/external/sqlite new file mode 160000 index 0000000000..f5fb820c0f --- /dev/null +++ b/external/sqlite @@ -0,0 +1 @@ +Subproject commit f5fb820c0f4781337faf02ed871be68d13a83d94 diff --git a/samples/api_buffered_tracing/client.cpp b/samples/api_buffered_tracing/client.cpp index cfe38f6149..721e0ac592 100644 --- a/samples/api_buffered_tracing/client.cpp +++ b/samples/api_buffered_tracing/client.cpp @@ -56,6 +56,7 @@ #include #include #include +#include #include #include #include @@ -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(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(user_data)->emplace_back( source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()}); diff --git a/samples/counter_collection/main.cpp b/samples/counter_collection/main.cpp index bdd4795a38..a33f52c815 100644 --- a/samples/counter_collection/main.cpp +++ b/samples/counter_collection/main.cpp @@ -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; } diff --git a/source/include/rocprofiler-sdk/callback_tracing.h b/source/include/rocprofiler-sdk/callback_tracing.h index 4776c9c6e4..8a6f648f2e 100644 --- a/source/include/rocprofiler-sdk/callback_tracing.h +++ b/source/include/rocprofiler-sdk/callback_tracing.h @@ -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; /** diff --git a/source/include/rocprofiler-sdk/cxx/CMakeLists.txt b/source/include/rocprofiler-sdk/cxx/CMakeLists.txt index a764cfd008..f4d1a93c1a 100644 --- a/source/include/rocprofiler-sdk/cxx/CMakeLists.txt +++ b/source/include/rocprofiler-sdk/cxx/CMakeLists.txt @@ -13,3 +13,4 @@ install( add_subdirectory(codeobj) add_subdirectory(details) +add_subdirectory(serialization) diff --git a/source/include/rocprofiler-sdk/cxx/details/mpl.hpp b/source/include/rocprofiler-sdk/cxx/details/mpl.hpp index ca89afcc69..5ace3a194a 100644 --- a/source/include/rocprofiler-sdk/cxx/details/mpl.hpp +++ b/source/include/rocprofiler-sdk/cxx/details/mpl.hpp @@ -249,6 +249,12 @@ reserve(ContainerT& _c, ArgT _arg) { return impl::reserve(_c, 0, _arg); } + +template +struct assert_false +{ + static constexpr auto value = false; +}; } // namespace mpl } // namespace sdk } // namespace rocprofiler diff --git a/source/include/rocprofiler-sdk/cxx/details/tokenize.hpp b/source/include/rocprofiler-sdk/cxx/details/tokenize.hpp index 40f9d19a94..9291afb374 100644 --- a/source/include/rocprofiler-sdk/cxx/details/tokenize.hpp +++ b/source/include/rocprofiler-sdk/cxx/details/tokenize.hpp @@ -28,6 +28,7 @@ #include #include #include +#include #include 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 diff --git a/source/include/rocprofiler-sdk/cxx/enum_string.hpp b/source/include/rocprofiler-sdk/cxx/enum_string.hpp index 6f550aa8f5..f3fd32a2bc 100644 --- a/source/include/rocprofiler-sdk/cxx/enum_string.hpp +++ b/source/include/rocprofiler-sdk/cxx/enum_string.hpp @@ -115,6 +115,7 @@ get_enum_label(EnumT val, std::index_sequence) return get_enum_label(val, std::index_sequence{}); return info::unsupported; + (void) val; // suppress unused-but-set-parameter warning } // clang-format off diff --git a/source/include/rocprofiler-sdk/cxx/perfetto.hpp b/source/include/rocprofiler-sdk/cxx/perfetto.hpp index 19ca2ca89e..f69f47615b 100644 --- a/source/include/rocprofiler-sdk/cxx/perfetto.hpp +++ b/source/include/rocprofiler-sdk/cxx/perfetto.hpp @@ -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 \ + { \ + 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 @@ -175,7 +193,101 @@ add_perfetto_annotation(perfetto_event_context_t& ctx, Np&& _name, Tp&& _val) } // namespace sdk } // namespace rocprofiler +#include + +namespace rocprofiler +{ +namespace sdk +{ +template +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 +const char* +get_perfetto_category(KindT kind, std::index_sequence) +{ + if(kind == Idx) return rocprofiler_tracing_perfetto_category::name; + // recursion until tail empty + if constexpr(sizeof...(Tail) > 0) + return get_perfetto_category(kind, std::index_sequence{}); + return nullptr; +} + +template +const char* +get_perfetto_category(KindT kind) +{ + if constexpr(std::is_same::value) + { + return get_perfetto_category( + kind, std::make_index_sequence{}); + } + else if constexpr(std::is_same::value) + { + return get_perfetto_category( + kind, std::make_index_sequence{}); + } + else + { + static_assert(rocprofiler::sdk::mpl::assert_false::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 diff --git a/source/include/rocprofiler-sdk/cxx/serialization.hpp b/source/include/rocprofiler-sdk/cxx/serialization.hpp index 97f2dae792..1492fbd6a7 100644 --- a/source/include/rocprofiler-sdk/cxx/serialization.hpp +++ b/source/include/rocprofiler-sdk/cxx/serialization.hpp @@ -23,1320 +23,5 @@ #pragma once -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include -#include -#include -#include - -#define ROCP_SDK_SAVE_DATA_FIELD(FIELD) ar(make_nvp(#FIELD, data.FIELD)) -#define ROCP_SDK_SAVE_DATA_VALUE(NAME, VALUE) ar(make_nvp(NAME, data.VALUE)) -#define ROCP_SDK_SAVE_VALUE(NAME, VALUE) ar(make_nvp(NAME, VALUE)) -#define ROCP_SDK_SAVE_DATA_CSTR(FIELD) \ - ar(make_nvp(#FIELD, std::string{data.FIELD ? data.FIELD : ""})) -#define ROCP_SDK_SAVE_DATA_BITFIELD(NAME, VALUE) \ - { \ - auto _val = data.VALUE; \ - ar(make_nvp(NAME, _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 - -namespace rocprofiler -{ -namespace sdk -{ -namespace serialization -{ -struct buffer_tracing_args -{ - std::string type = {}; - std::string name = {}; - std::string value = {}; -}; - -template -auto -get_buffer_tracing_args(Tp& data) -{ - auto populate_args_array = [](rocprofiler_buffer_tracing_kind_t /*kind*/, - rocprofiler_tracing_operation_t /*operation*/, - uint32_t arg_number, - const void* const /*arg_value_addr*/, - int32_t /*arg_indirection_count*/, - const char* arg_type, - const char* arg_name, - const char* arg_value_str, - void* cb_data) -> int { - if(!cb_data) return 1; - - auto* vec = static_cast*>(cb_data); - auto sz = std::max(arg_number + 1, vec->size()); - vec->resize(sz, buffer_tracing_args{}); - vec->at(arg_number) = buffer_tracing_args{arg_type, arg_name, arg_value_str}; - - return 0; - }; - - auto ret = std::vector{}; - auto record = rocprofiler_record_header_t{}; - record.hash = - rocprofiler_record_header_compute_hash(ROCPROFILER_BUFFER_CATEGORY_TRACING, data.kind); - record.payload = &data; - - rocprofiler_iterate_buffer_tracing_record_args(record, populate_args_array, &ret); - - return ret; -} -} // namespace serialization -} // namespace sdk -} // namespace rocprofiler - -ROCPROFILER_SDK_CEREAL_NAMESPACE_BEGIN - -namespace sdk = ::rocprofiler::sdk; - -template -void -save(ArchiveT& ar, const sdk::serialization::buffer_tracing_args& data) -{ - ROCP_SDK_SAVE_DATA_FIELD(type); - ROCP_SDK_SAVE_DATA_FIELD(name); - ROCP_SDK_SAVE_DATA_FIELD(value); -} - -template -void -save(ArchiveT& ar, rocprofiler_context_id_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(handle); -} - -template -void -save(ArchiveT& ar, rocprofiler_agent_id_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(handle); -} - -template -void -save(ArchiveT& ar, hsa_agent_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(handle); -} - -template -void -save(ArchiveT& ar, rocprofiler_queue_id_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(handle); -} - -template -void -save(ArchiveT& ar, rocprofiler_stream_id_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(handle); -} - -template -void -save(ArchiveT& ar, rocprofiler_counter_id_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(handle); -} - -template -void -save(ArchiveT& ar, rocprofiler_correlation_id_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(internal); - ROCP_SDK_SAVE_DATA_VALUE("external", external.value); - ROCP_SDK_SAVE_DATA_VALUE("ancestor", ancestor); -} - -template -void -save(ArchiveT& ar, rocprofiler_async_correlation_id_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(internal); - ROCP_SDK_SAVE_DATA_VALUE("external", external.value); -} - -template -void -save(ArchiveT& ar, rocprofiler_dim3_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(x); - ROCP_SDK_SAVE_DATA_FIELD(y); - ROCP_SDK_SAVE_DATA_FIELD(z); -} - -template -void -save(ArchiveT& ar, rocprofiler_address_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(handle); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_code_object_load_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(code_object_id); - ROCP_SDK_SAVE_DATA_FIELD(agent_id); - ROCP_SDK_SAVE_DATA_CSTR(uri); - ROCP_SDK_SAVE_DATA_FIELD(load_base); - ROCP_SDK_SAVE_DATA_FIELD(load_size); - ROCP_SDK_SAVE_DATA_FIELD(load_delta); - ROCP_SDK_SAVE_DATA_FIELD(storage_type); - if(data.storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_FILE) - { - ROCP_SDK_SAVE_DATA_FIELD(storage_file); - } - else if(data.storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_MEMORY) - { - ROCP_SDK_SAVE_DATA_FIELD(memory_base); - ROCP_SDK_SAVE_DATA_FIELD(memory_size); - } -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(kernel_id); - ROCP_SDK_SAVE_DATA_FIELD(code_object_id); - ROCP_SDK_SAVE_DATA_CSTR(kernel_name); - ROCP_SDK_SAVE_DATA_FIELD(kernel_object); - ROCP_SDK_SAVE_DATA_FIELD(kernarg_segment_size); - ROCP_SDK_SAVE_DATA_FIELD(kernarg_segment_alignment); - ROCP_SDK_SAVE_DATA_FIELD(group_segment_size); - ROCP_SDK_SAVE_DATA_FIELD(private_segment_size); - ROCP_SDK_SAVE_DATA_FIELD(sgpr_count); - ROCP_SDK_SAVE_DATA_FIELD(arch_vgpr_count); - ROCP_SDK_SAVE_DATA_FIELD(accum_vgpr_count); - ROCP_SDK_SAVE_DATA_FIELD(kernel_code_entry_byte_offset); - ROCP_SDK_SAVE_DATA_FIELD(kernel_address); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_code_object_host_kernel_symbol_register_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(host_function_id); - ROCP_SDK_SAVE_DATA_FIELD(kernel_id); - ROCP_SDK_SAVE_DATA_FIELD(code_object_id); - ROCP_SDK_SAVE_DATA_FIELD(host_function); - ROCP_SDK_SAVE_DATA_FIELD(modules); - ROCP_SDK_SAVE_DATA_CSTR(device_function); - ROCP_SDK_SAVE_DATA_FIELD(thread_limit); - ROCP_SDK_SAVE_DATA_FIELD(thread_ids); - ROCP_SDK_SAVE_DATA_FIELD(block_ids); - ROCP_SDK_SAVE_DATA_FIELD(block_dims); - ROCP_SDK_SAVE_DATA_FIELD(grid_dims); - ROCP_SDK_SAVE_DATA_FIELD(workgroup_size); -} - -template -void -save(ArchiveT& ar, rocprofiler_hsa_api_retval_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(uint64_t_retval); -} - -template -void -save(ArchiveT& ar, const hsa_queue_t& data) -{ - ar(make_nvp("queue_id", data.id)); -} - -template -void -save(ArchiveT& ar, hsa_amd_event_scratch_alloc_start_t data) -{ - ar(make_nvp("queue_id", *data.queue)); - ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); -} - -template -void -save(ArchiveT& ar, hsa_amd_event_scratch_alloc_end_t data) -{ - ar(make_nvp("queue_id", *data.queue)); - ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(num_slots); - ROCP_SDK_SAVE_DATA_FIELD(flags); -} - -template -void -save(ArchiveT& ar, hsa_amd_event_scratch_free_start_t data) -{ - ar(make_nvp("queue_id", *data.queue)); -} - -template -void -save(ArchiveT& ar, hsa_amd_event_scratch_free_end_t data) -{ - ar(make_nvp("queue_id", *data.queue)); - ROCP_SDK_SAVE_DATA_FIELD(flags); -} - -template -void -save(ArchiveT& ar, hsa_amd_event_scratch_async_reclaim_start_t data) -{ - ar(make_nvp("queue_id", *data.queue)); -} - -template -void -save(ArchiveT& ar, hsa_amd_event_scratch_async_reclaim_end_t data) -{ - ar(make_nvp("queue_id", *data.queue)); - ROCP_SDK_SAVE_DATA_FIELD(flags); -} - -template -void -save(ArchiveT& ar, rocprofiler_marker_api_retval_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(int64_t_retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_hsa_api_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - // ROCP_SDK_SAVE_DATA_FIELD(args); - ROCP_SDK_SAVE_DATA_FIELD(retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_marker_api_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - // ROCP_SDK_SAVE_DATA_FIELD(args); - ROCP_SDK_SAVE_DATA_FIELD(retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_hip_api_retval_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(uint64_t_retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_hip_api_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - // ROCP_SDK_SAVE_DATA_FIELD(args); - ROCP_SDK_SAVE_DATA_FIELD(retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_scratch_memory_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(agent_id); - ROCP_SDK_SAVE_DATA_FIELD(queue_id); - ROCP_SDK_SAVE_DATA_FIELD(flags); - ROCP_SDK_SAVE_DATA_FIELD(args_kind); -} - -template -void -save(ArchiveT& ar, rocprofiler_kernel_dispatch_info_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(agent_id); - ROCP_SDK_SAVE_DATA_FIELD(queue_id); - ROCP_SDK_SAVE_DATA_FIELD(kernel_id); - ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); - ROCP_SDK_SAVE_DATA_FIELD(private_segment_size); - ROCP_SDK_SAVE_DATA_FIELD(group_segment_size); - ROCP_SDK_SAVE_DATA_FIELD(workgroup_size); - ROCP_SDK_SAVE_DATA_FIELD(grid_size); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_kernel_dispatch_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(dispatch_info); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_memory_copy_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(dst_agent_id); - ROCP_SDK_SAVE_DATA_FIELD(src_agent_id); - ROCP_SDK_SAVE_DATA_FIELD(bytes); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_memory_allocation_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(agent_id); - ROCP_SDK_SAVE_VALUE("address", rocprofiler::sdk::utility::as_hex(data.address.handle, 16)); - ROCP_SDK_SAVE_DATA_FIELD(allocation_size); -} - -template -void -save(ArchiveT& ar, rocprofiler_rccl_api_retval_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(ncclResult_t_retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_rccl_api_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - // ROCP_SDK_SAVE_DATA_FIELD(args); - ROCP_SDK_SAVE_DATA_FIELD(retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_rocdecode_api_retval_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(uint64_t_retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_rocdecode_api_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_rocjpeg_api_retval_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(rocJpegStatus_retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_rocjpeg_api_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_ompt_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - // ROCP_SDK_SAVE_DATA_FIELD(args); -} - -template -void -save(ArchiveT& ar, rocprofiler_dispatch_counting_service_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(correlation_id); - ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(dispatch_info); -} - -template -void -save(ArchiveT& ar, rocprofiler_dispatch_counting_service_record_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(num_records); - ROCP_SDK_SAVE_DATA_FIELD(correlation_id); - ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(dispatch_info); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_record_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(context_id); - ROCP_SDK_SAVE_DATA_FIELD(thread_id); - ROCP_SDK_SAVE_DATA_FIELD(kind); - ROCP_SDK_SAVE_DATA_FIELD(operation); - ROCP_SDK_SAVE_DATA_FIELD(correlation_id); - ROCP_SDK_SAVE_DATA_FIELD(phase); -} - -template -void -save_buffer_tracing_api_record(ArchiveT& ar, Tp data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(kind); - ROCP_SDK_SAVE_DATA_FIELD(operation); - ROCP_SDK_SAVE_DATA_FIELD(correlation_id); - ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(thread_id); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_hsa_api_record_t data) -{ - save_buffer_tracing_api_record(ar, data); -} - -template -void -save(ArchiveT& ar, rocprofiler_counter_record_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(id); - ROCP_SDK_SAVE_DATA_FIELD(counter_value); - ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_hip_api_record_t data) -{ - save_buffer_tracing_api_record(ar, data); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_hip_api_ext_record_t data) -{ - save_buffer_tracing_api_record(ar, data); - auto args = sdk::serialization::get_buffer_tracing_args(data); - ROCP_SDK_SAVE_VALUE("args", args); - ROCP_SDK_SAVE_DATA_FIELD(retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_marker_api_record_t data) -{ - save_buffer_tracing_api_record(ar, data); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_rccl_api_record_t data) -{ - save_buffer_tracing_api_record(ar, data); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_rocdecode_api_record_t data) -{ - save_buffer_tracing_api_record(ar, data); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_rocdecode_api_ext_record_t data) -{ - save_buffer_tracing_api_record(ar, data); - auto args = sdk::serialization::get_buffer_tracing_args(data); - ROCP_SDK_SAVE_VALUE("args", args); - ROCP_SDK_SAVE_DATA_FIELD(retval); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_rocjpeg_api_record_t data) -{ - save_buffer_tracing_api_record(ar, data); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_ompt_target_t data) -{ - ROCP_SDK_SAVE_DATA_VALUE("kind", kind); - ROCP_SDK_SAVE_DATA_VALUE("device", device_num); - ROCP_SDK_SAVE_DATA_VALUE("task_id", task_id); - ROCP_SDK_SAVE_DATA_VALUE("target_id", target_id); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_ompt_target_data_op_t data) -{ - ROCP_SDK_SAVE_DATA_VALUE("host_op_id", host_op_id); - ROCP_SDK_SAVE_DATA_VALUE("optype", optype); - ROCP_SDK_SAVE_DATA_VALUE("src_device_num", src_device_num); - ROCP_SDK_SAVE_DATA_VALUE("dst_device_num", dst_device_num); - ROCP_SDK_SAVE_DATA_VALUE("bytes", bytes); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_ompt_target_kernel_t data) -{ - ROCP_SDK_SAVE_DATA_VALUE("host_op_id", host_op_id); - ROCP_SDK_SAVE_DATA_VALUE("device_num", device_num); - ROCP_SDK_SAVE_DATA_VALUE("requested_num_teams", requested_num_teams); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_ompt_record_t data) -{ - save_buffer_tracing_api_record(ar, data); - - if(data.operation == ROCPROFILER_OMPT_ID_target_emi) - { - ROCP_SDK_SAVE_DATA_FIELD(target); - } - else if(data.operation == ROCPROFILER_OMPT_ID_target_data_op_emi) - { - ROCP_SDK_SAVE_DATA_FIELD(target_data_op); - } - else if(data.operation == ROCPROFILER_OMPT_ID_target_submit_emi) - { - ROCP_SDK_SAVE_DATA_FIELD(target_kernel); - } -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_kernel_dispatch_record_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(kind); - ROCP_SDK_SAVE_DATA_FIELD(operation); - ROCP_SDK_SAVE_DATA_FIELD(thread_id); - ROCP_SDK_SAVE_DATA_FIELD(correlation_id); - ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(dispatch_info); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_memory_copy_record_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(kind); - ROCP_SDK_SAVE_DATA_FIELD(operation); - ROCP_SDK_SAVE_DATA_FIELD(thread_id); - ROCP_SDK_SAVE_DATA_FIELD(correlation_id); - ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(dst_agent_id); - ROCP_SDK_SAVE_DATA_FIELD(src_agent_id); - ROCP_SDK_SAVE_DATA_FIELD(bytes); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_memory_allocation_record_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(kind); - ROCP_SDK_SAVE_DATA_FIELD(operation); - ROCP_SDK_SAVE_DATA_FIELD(thread_id); - ROCP_SDK_SAVE_DATA_FIELD(correlation_id); - ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(agent_id); - ROCP_SDK_SAVE_VALUE("address", rocprofiler::sdk::utility::as_hex(data.address.handle, 16)); - ROCP_SDK_SAVE_DATA_FIELD(allocation_size); -} - -template -void -save(ArchiveT& ar, const rocprofiler_page_migration_page_fault_start_t& data) -{ - ROCP_SDK_SAVE_DATA_BITFIELD("read_fault", read_fault); - ROCP_SDK_SAVE_DATA_FIELD(agent_id); - ROCP_SDK_SAVE_DATA_FIELD(address); -} - -template -void -save(ArchiveT& ar, const rocprofiler_page_migration_page_fault_end_t& data) -{ - ROCP_SDK_SAVE_DATA_BITFIELD("migrated", migrated); - ROCP_SDK_SAVE_DATA_FIELD(agent_id); - ROCP_SDK_SAVE_DATA_FIELD(address); -} - -template -void -save(ArchiveT& ar, const rocprofiler_page_migration_page_migrate_start_t& data) -{ - ROCP_SDK_SAVE_DATA_FIELD(start_addr); - ROCP_SDK_SAVE_DATA_FIELD(end_addr); - ROCP_SDK_SAVE_DATA_FIELD(from_agent); - ROCP_SDK_SAVE_DATA_FIELD(to_agent); - ROCP_SDK_SAVE_DATA_FIELD(prefetch_agent); - ROCP_SDK_SAVE_DATA_FIELD(preferred_agent); - ROCP_SDK_SAVE_DATA_FIELD(trigger); -} - -template -void -save(ArchiveT& ar, const rocprofiler_page_migration_page_migrate_end_t& data) -{ - ROCP_SDK_SAVE_DATA_FIELD(start_addr); - ROCP_SDK_SAVE_DATA_FIELD(end_addr); - ROCP_SDK_SAVE_DATA_FIELD(from_agent); - ROCP_SDK_SAVE_DATA_FIELD(to_agent); - ROCP_SDK_SAVE_DATA_FIELD(trigger); - ROCP_SDK_SAVE_DATA_FIELD(error_code); -} - -template -void -save(ArchiveT& ar, const rocprofiler_page_migration_queue_eviction_t& data) -{ - ROCP_SDK_SAVE_DATA_FIELD(agent_id); - ROCP_SDK_SAVE_DATA_FIELD(trigger); -} - -template -void -save(ArchiveT& ar, const rocprofiler_page_migration_queue_restore_t& data) -{ - ROCP_SDK_SAVE_DATA_BITFIELD("rescheduled", rescheduled); - ROCP_SDK_SAVE_DATA_FIELD(agent_id); -} - -template -void -save(ArchiveT& ar, const rocprofiler_page_migration_unmap_from_gpu_t& data) -{ - ROCP_SDK_SAVE_DATA_FIELD(start_addr); - ROCP_SDK_SAVE_DATA_FIELD(end_addr); - ROCP_SDK_SAVE_DATA_FIELD(agent_id); - ROCP_SDK_SAVE_DATA_FIELD(trigger); -} - -template -void -save(ArchiveT& ar, const rocprofiler_page_migration_dropped_event_t& data) -{ - ROCP_SDK_SAVE_DATA_FIELD(dropped_events_count); -} - -namespace details -{ -template -struct save_page_migration_arg; - -#define ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(ENUM_VALUE, UNION_ARG) \ - template <> \ - struct save_page_migration_arg \ - { \ - static constexpr auto value = ROCPROFILER_PAGE_MIGRATION_##ENUM_VALUE; \ - template \ - void operator()(ArchiveT& ar, rocprofiler_page_migration_args_t args) \ - { \ - ar(make_nvp(#UNION_ARG, args.UNION_ARG)); \ - } \ - }; - -ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(NONE, none) -ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(PAGE_MIGRATE_START, page_migrate_start) -ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(PAGE_MIGRATE_END, page_migrate_end) -ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(PAGE_FAULT_START, page_fault_start) -ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(PAGE_FAULT_END, page_fault_end) -ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(QUEUE_EVICTION, queue_eviction) -ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(QUEUE_RESTORE, queue_restore) -ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(UNMAP_FROM_GPU, unmap_from_gpu) -ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(DROPPED_EVENT, dropped_event) - -#undef ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG - -template -void -save_page_migration_args(ArchiveT& ar, - rocprofiler_page_migration_operation_t op, - rocprofiler_page_migration_args_t args, - std::index_sequence) -{ - using save_page_migration_type = save_page_migration_arg; - if(op == save_page_migration_type::value) - { - if constexpr(save_page_migration_type::value != ROCPROFILER_PAGE_MIGRATION_NONE) - save_page_migration_type{}(ar, args); - } - else if constexpr(sizeof...(IdxTail) > 0) - { - save_page_migration_args(ar, op, args, std::index_sequence{}); - } -} -} // namespace details - -template -void -save(ArchiveT& ar, const rocprofiler_buffer_tracing_page_migration_record_t& data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(kind); - ROCP_SDK_SAVE_DATA_FIELD(operation); - ROCP_SDK_SAVE_DATA_FIELD(timestamp); - ROCP_SDK_SAVE_DATA_FIELD(pid); - details::save_page_migration_args( - ar, data.operation, data.args, std::make_index_sequence{}); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_scratch_memory_record_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(kind); - ROCP_SDK_SAVE_DATA_FIELD(operation); - ROCP_SDK_SAVE_DATA_FIELD(agent_id); - ROCP_SDK_SAVE_DATA_FIELD(queue_id); - ROCP_SDK_SAVE_DATA_FIELD(thread_id); - ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); - ROCP_SDK_SAVE_DATA_FIELD(correlation_id); - ROCP_SDK_SAVE_DATA_FIELD(flags); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_correlation_id_retirement_record_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(kind); - ROCP_SDK_SAVE_DATA_FIELD(timestamp); - ROCP_SDK_SAVE_DATA_FIELD(internal_correlation_id); -} - -template -void -save(ArchiveT& ar, HsaCacheType data) -{ - ROCP_SDK_SAVE_DATA_BITFIELD("Data", ui32.Data); - ROCP_SDK_SAVE_DATA_BITFIELD("Instruction", ui32.Instruction); - ROCP_SDK_SAVE_DATA_BITFIELD("CPU", ui32.CPU); - ROCP_SDK_SAVE_DATA_BITFIELD("HSACU", ui32.HSACU); -} - -template -void -save(ArchiveT& ar, HSA_LINKPROPERTY data) -{ - ROCP_SDK_SAVE_DATA_BITFIELD("Override", ui32.Override); - ROCP_SDK_SAVE_DATA_BITFIELD("NonCoherent", ui32.NonCoherent); - ROCP_SDK_SAVE_DATA_BITFIELD("NoAtomics32bit", ui32.NoAtomics32bit); - ROCP_SDK_SAVE_DATA_BITFIELD("NoAtomics64bit", ui32.NoAtomics64bit); - ROCP_SDK_SAVE_DATA_BITFIELD("NoPeerToPeerDMA", ui32.NoPeerToPeerDMA); -} - -template -void -save(ArchiveT& ar, HSA_CAPABILITY data) -{ - ROCP_SDK_SAVE_DATA_BITFIELD("HotPluggable", ui32.HotPluggable); - ROCP_SDK_SAVE_DATA_BITFIELD("HSAMMUPresent", ui32.HSAMMUPresent); - ROCP_SDK_SAVE_DATA_BITFIELD("SharedWithGraphics", ui32.SharedWithGraphics); - ROCP_SDK_SAVE_DATA_BITFIELD("QueueSizePowerOfTwo", ui32.QueueSizePowerOfTwo); - ROCP_SDK_SAVE_DATA_BITFIELD("QueueSize32bit", ui32.QueueSize32bit); - ROCP_SDK_SAVE_DATA_BITFIELD("QueueIdleEvent", ui32.QueueIdleEvent); - ROCP_SDK_SAVE_DATA_BITFIELD("VALimit", ui32.VALimit); - ROCP_SDK_SAVE_DATA_BITFIELD("WatchPointsSupported", ui32.WatchPointsSupported); - ROCP_SDK_SAVE_DATA_BITFIELD("WatchPointsTotalBits", ui32.WatchPointsTotalBits); - ROCP_SDK_SAVE_DATA_BITFIELD("DoorbellType", ui32.DoorbellType); - ROCP_SDK_SAVE_DATA_BITFIELD("AQLQueueDoubleMap", ui32.AQLQueueDoubleMap); - ROCP_SDK_SAVE_DATA_BITFIELD("DebugTrapSupported", ui32.DebugTrapSupported); - ROCP_SDK_SAVE_DATA_BITFIELD("WaveLaunchTrapOverrideSupported", - ui32.WaveLaunchTrapOverrideSupported); - ROCP_SDK_SAVE_DATA_BITFIELD("WaveLaunchModeSupported", ui32.WaveLaunchModeSupported); - ROCP_SDK_SAVE_DATA_BITFIELD("PreciseMemoryOperationsSupported", - ui32.PreciseMemoryOperationsSupported); - ROCP_SDK_SAVE_DATA_BITFIELD("DEPRECATED_SRAM_EDCSupport", ui32.DEPRECATED_SRAM_EDCSupport); - ROCP_SDK_SAVE_DATA_BITFIELD("Mem_EDCSupport", ui32.Mem_EDCSupport); - ROCP_SDK_SAVE_DATA_BITFIELD("RASEventNotify", ui32.RASEventNotify); - ROCP_SDK_SAVE_DATA_BITFIELD("ASICRevision", ui32.ASICRevision); - ROCP_SDK_SAVE_DATA_BITFIELD("SRAM_EDCSupport", ui32.SRAM_EDCSupport); - ROCP_SDK_SAVE_DATA_BITFIELD("SVMAPISupported", ui32.SVMAPISupported); - ROCP_SDK_SAVE_DATA_BITFIELD("CoherentHostAccess", ui32.CoherentHostAccess); - ROCP_SDK_SAVE_DATA_BITFIELD("DebugSupportedFirmware", ui32.DebugSupportedFirmware); -} - -template -void -save(ArchiveT& ar, HSA_MEMORYPROPERTY data) -{ - ROCP_SDK_SAVE_DATA_BITFIELD("HotPluggable", ui32.HotPluggable); - ROCP_SDK_SAVE_DATA_BITFIELD("NonVolatile", ui32.NonVolatile); -} - -template -void -save(ArchiveT& ar, HSA_ENGINE_VERSION data) -{ - ROCP_SDK_SAVE_DATA_BITFIELD("uCodeSDMA", uCodeSDMA); - ROCP_SDK_SAVE_DATA_BITFIELD("uCodeRes", uCodeRes); -} - -template -void -save(ArchiveT& ar, HSA_ENGINE_ID data) -{ - ROCP_SDK_SAVE_DATA_BITFIELD("uCode", ui32.uCode); - ROCP_SDK_SAVE_DATA_BITFIELD("Major", ui32.Major); - ROCP_SDK_SAVE_DATA_BITFIELD("Minor", ui32.Minor); - ROCP_SDK_SAVE_DATA_BITFIELD("Stepping", ui32.Stepping); -} - -template -void -save(ArchiveT& ar, rocprofiler_agent_cache_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(processor_id_low); - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(level); - ROCP_SDK_SAVE_DATA_FIELD(cache_line_size); - ROCP_SDK_SAVE_DATA_FIELD(cache_lines_per_tag); - ROCP_SDK_SAVE_DATA_FIELD(association); - ROCP_SDK_SAVE_DATA_FIELD(latency); - ROCP_SDK_SAVE_DATA_FIELD(type); -} -template -void -save(ArchiveT& ar, rocprofiler_pc_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(code_object_id); - ROCP_SDK_SAVE_DATA_FIELD(code_object_offset); -} - -template -void -save(ArchiveT& ar, rocprofiler_pc_sampling_hw_id_v0_t data) -{ - ROCP_SDK_SAVE_DATA_BITFIELD("chiplet", chiplet); - ROCP_SDK_SAVE_DATA_BITFIELD("wave_id", wave_id); - ROCP_SDK_SAVE_DATA_BITFIELD("simd_id", simd_id); - ROCP_SDK_SAVE_DATA_BITFIELD("pipe_id", pipe_id); - ROCP_SDK_SAVE_DATA_BITFIELD("cu_or_wgp_id", cu_or_wgp_id); - ROCP_SDK_SAVE_DATA_BITFIELD("shader_array_id", shader_array_id); - ROCP_SDK_SAVE_DATA_BITFIELD("shader_engine_id", shader_engine_id); - ROCP_SDK_SAVE_DATA_BITFIELD("workgroup_id ", workgroup_id); - ROCP_SDK_SAVE_DATA_BITFIELD("vm_id", vm_id); - ROCP_SDK_SAVE_DATA_BITFIELD("queue_id", queue_id); - ROCP_SDK_SAVE_DATA_BITFIELD("microengine_id", microengine_id); -} - -template -void -save(ArchiveT& ar, rocprofiler_pc_sampling_record_host_trap_v0_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(hw_id); - ROCP_SDK_SAVE_DATA_FIELD(pc); - ROCP_SDK_SAVE_DATA_FIELD(exec_mask); - ROCP_SDK_SAVE_DATA_FIELD(timestamp); - ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); - ROCP_SDK_SAVE_DATA_VALUE("corr_id", correlation_id); - ROCP_SDK_SAVE_DATA_VALUE("wrkgrp_id", workgroup_id); - ROCP_SDK_SAVE_DATA_BITFIELD("wave_in_grp", wave_in_group); -} - -template -void -save(ArchiveT& ar, rocprofiler_pc_sampling_record_stochastic_header_t data) -{ - ROCP_SDK_SAVE_DATA_BITFIELD("has_mem_cnt", has_memory_counter); -} - -template -void -save_pc_sampling_inst_type(ArchiveT& ar, rocprofiler_pc_sampling_instruction_type_t inst_type) -{ - ar(make_nvp("inst_type", - std::string(rocprofiler_get_pc_sampling_instruction_type_name(inst_type)))); -} - -template -void -save_pc_sampling_stall_reason(ArchiveT& ar, - rocprofiler_pc_sampling_instruction_not_issued_reason_t stall_reason) -{ - ar(make_nvp( - "stall_reason", - std::string(rocprofiler_get_pc_sampling_instruction_not_issued_reason_name(stall_reason)))); -} - -template -void -save(ArchiveT& ar, rocprofiler_pc_sampling_snapshot_v0_t data) -{ - save_pc_sampling_stall_reason( - ar, - static_cast( - data.reason_not_issued)); - - ROCP_SDK_SAVE_DATA_BITFIELD("dual_issue_valu", dual_issue_valu); - - // Arb state (pipe issued) - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_valu", arb_state_issue_valu); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_matrix", arb_state_issue_matrix); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_lds", arb_state_issue_lds); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_lds_direct", arb_state_issue_lds_direct); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_scalar", arb_state_issue_scalar); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_vmem_tex", arb_state_issue_vmem_tex); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_flat", arb_state_issue_flat); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_exp", arb_state_issue_exp); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_misc", arb_state_issue_misc); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_brmsg", arb_state_issue_brmsg); - // Arb state (pipe stalled) - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_valu", arb_state_stall_valu); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_matrix", arb_state_stall_matrix); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_lds", arb_state_stall_lds); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_lds_direct", arb_state_stall_lds_direct); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_scalar", arb_state_stall_scalar); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_vmem_tex", arb_state_stall_vmem_tex); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_flat", arb_state_stall_flat); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_exp", arb_state_stall_exp); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_misc", arb_state_stall_misc); - ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_brmsg", arb_state_stall_brmsg); -} - -template -void -save(ArchiveT& ar, rocprofiler_pc_sampling_record_stochastic_v0_t data) -{ - // flags specific for stochastic sampling - ROCP_SDK_SAVE_DATA_FIELD(flags); - - // Common for host-trap and stochastic - ROCP_SDK_SAVE_DATA_FIELD(hw_id); - ROCP_SDK_SAVE_DATA_FIELD(pc); - ROCP_SDK_SAVE_DATA_FIELD(exec_mask); - ROCP_SDK_SAVE_DATA_FIELD(timestamp); - ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); - ROCP_SDK_SAVE_DATA_VALUE("corr_id", correlation_id); - ROCP_SDK_SAVE_DATA_VALUE("wrkgrp_id", workgroup_id); - ROCP_SDK_SAVE_DATA_BITFIELD("wave_in_grp", wave_in_group); - - // fields specific for stochastic - ROCP_SDK_SAVE_DATA_BITFIELD("wave_issued", wave_issued); - save_pc_sampling_inst_type( - ar, static_cast(data.inst_type)); - ROCP_SDK_SAVE_DATA_BITFIELD("wave_cnt", wave_count); - ROCP_SDK_SAVE_DATA_FIELD(snapshot); - - // TODO: add memory counters -} - -template -void -save(ArchiveT& ar, rocprofiler_agent_io_link_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(type); - ROCP_SDK_SAVE_DATA_FIELD(version_major); - ROCP_SDK_SAVE_DATA_FIELD(version_minor); - ROCP_SDK_SAVE_DATA_FIELD(node_from); - ROCP_SDK_SAVE_DATA_FIELD(node_to); - ROCP_SDK_SAVE_DATA_FIELD(weight); - ROCP_SDK_SAVE_DATA_FIELD(min_latency); - ROCP_SDK_SAVE_DATA_FIELD(max_latency); - ROCP_SDK_SAVE_DATA_FIELD(min_bandwidth); - ROCP_SDK_SAVE_DATA_FIELD(max_bandwidth); - ROCP_SDK_SAVE_DATA_FIELD(recommended_transfer_size); - ROCP_SDK_SAVE_DATA_FIELD(flags); -} - -template -void -save(ArchiveT& ar, rocprofiler_agent_mem_bank_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(heap_type); - ROCP_SDK_SAVE_DATA_FIELD(flags); - ROCP_SDK_SAVE_DATA_FIELD(width); - ROCP_SDK_SAVE_DATA_FIELD(mem_clk_max); - ROCP_SDK_SAVE_DATA_FIELD(size_in_bytes); -} - -template -void -save(ArchiveT& ar, rocprofiler_pc_sampling_configuration_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(method); - ROCP_SDK_SAVE_DATA_FIELD(unit); - ROCP_SDK_SAVE_DATA_FIELD(min_interval); - ROCP_SDK_SAVE_DATA_FIELD(max_interval); - ROCP_SDK_SAVE_DATA_FIELD(flags); -} - -template -void -save(ArchiveT& ar, rocprofiler_agent_runtime_visiblity_t data) -{ - ROCP_SDK_SAVE_DATA_BITFIELD("hsa", hsa); - ROCP_SDK_SAVE_DATA_BITFIELD("hip", hip); - ROCP_SDK_SAVE_DATA_BITFIELD("rccl", rccl); - ROCP_SDK_SAVE_DATA_BITFIELD("rocdecode", rocdecode); -} - -template -void -save(ArchiveT& ar, rocprofiler_uuid_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(value); -} - -template -void -save(ArchiveT& ar, const rocprofiler_agent_v0_t& data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(id); - ROCP_SDK_SAVE_DATA_FIELD(type); - ROCP_SDK_SAVE_DATA_FIELD(cpu_cores_count); - ROCP_SDK_SAVE_DATA_FIELD(simd_count); - ROCP_SDK_SAVE_DATA_FIELD(mem_banks_count); - ROCP_SDK_SAVE_DATA_FIELD(caches_count); - ROCP_SDK_SAVE_DATA_FIELD(io_links_count); - ROCP_SDK_SAVE_DATA_FIELD(cpu_core_id_base); - ROCP_SDK_SAVE_DATA_FIELD(simd_id_base); - ROCP_SDK_SAVE_DATA_FIELD(max_waves_per_simd); - ROCP_SDK_SAVE_DATA_FIELD(lds_size_in_kb); - ROCP_SDK_SAVE_DATA_FIELD(gds_size_in_kb); - ROCP_SDK_SAVE_DATA_FIELD(num_gws); - ROCP_SDK_SAVE_DATA_FIELD(wave_front_size); - ROCP_SDK_SAVE_DATA_FIELD(num_xcc); - ROCP_SDK_SAVE_DATA_FIELD(cu_count); - ROCP_SDK_SAVE_DATA_FIELD(array_count); - ROCP_SDK_SAVE_DATA_FIELD(num_shader_banks); - ROCP_SDK_SAVE_DATA_FIELD(simd_arrays_per_engine); - ROCP_SDK_SAVE_DATA_FIELD(cu_per_simd_array); - ROCP_SDK_SAVE_DATA_FIELD(simd_per_cu); - ROCP_SDK_SAVE_DATA_FIELD(max_slots_scratch_cu); - ROCP_SDK_SAVE_DATA_FIELD(gfx_target_version); - ROCP_SDK_SAVE_DATA_FIELD(vendor_id); - ROCP_SDK_SAVE_DATA_FIELD(device_id); - ROCP_SDK_SAVE_DATA_FIELD(location_id); - ROCP_SDK_SAVE_DATA_FIELD(domain); - ROCP_SDK_SAVE_DATA_FIELD(drm_render_minor); - ROCP_SDK_SAVE_DATA_FIELD(num_sdma_engines); - ROCP_SDK_SAVE_DATA_FIELD(num_sdma_xgmi_engines); - ROCP_SDK_SAVE_DATA_FIELD(num_sdma_queues_per_engine); - ROCP_SDK_SAVE_DATA_FIELD(num_cp_queues); - ROCP_SDK_SAVE_DATA_FIELD(max_engine_clk_ccompute); - ROCP_SDK_SAVE_DATA_FIELD(max_engine_clk_fcompute); - ROCP_SDK_SAVE_DATA_FIELD(sdma_fw_version); - ROCP_SDK_SAVE_DATA_FIELD(fw_version); - ROCP_SDK_SAVE_DATA_FIELD(capability); - ROCP_SDK_SAVE_DATA_FIELD(cu_per_engine); - ROCP_SDK_SAVE_DATA_FIELD(max_waves_per_cu); - ROCP_SDK_SAVE_DATA_FIELD(family_id); - ROCP_SDK_SAVE_DATA_FIELD(workgroup_max_size); - ROCP_SDK_SAVE_DATA_FIELD(grid_max_size); - ROCP_SDK_SAVE_DATA_FIELD(local_mem_size); - ROCP_SDK_SAVE_DATA_FIELD(hive_id); - ROCP_SDK_SAVE_DATA_FIELD(gpu_id); - ROCP_SDK_SAVE_DATA_FIELD(workgroup_max_dim); - ROCP_SDK_SAVE_DATA_FIELD(grid_max_dim); - ROCP_SDK_SAVE_DATA_CSTR(name); - ROCP_SDK_SAVE_DATA_CSTR(vendor_name); - ROCP_SDK_SAVE_DATA_CSTR(product_name); - ROCP_SDK_SAVE_DATA_CSTR(model_name); - ROCP_SDK_SAVE_DATA_FIELD(node_id); - ROCP_SDK_SAVE_DATA_FIELD(logical_node_id); - ROCP_SDK_SAVE_DATA_FIELD(logical_node_type_id); - ROCP_SDK_SAVE_DATA_FIELD(runtime_visibility); - ROCP_SDK_SAVE_DATA_FIELD(uuid); - - auto generate = [&](auto name, const auto* value, uint64_t size) { - using value_type = std::remove_const_t>; - auto vec = std::vector{}; - vec.reserve(size); - for(uint64_t i = 0; i < size; ++i) - vec.emplace_back(value[i]); - ar(make_nvp(name, vec)); - }; - - 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 -void -save(ArchiveT& ar, rocprofiler_counter_info_v0_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(id); - ROCP_SDK_SAVE_DATA_BITFIELD("is_constant", is_constant); - ROCP_SDK_SAVE_DATA_BITFIELD("is_derived", is_derived); - ROCP_SDK_SAVE_DATA_CSTR(name); - ROCP_SDK_SAVE_DATA_CSTR(description); - ROCP_SDK_SAVE_DATA_CSTR(block); - ROCP_SDK_SAVE_DATA_CSTR(expression); -} - -template -void -save(ArchiveT& ar, rocprofiler_counter_info_v1_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(id); - ROCP_SDK_SAVE_DATA_BITFIELD("is_constant", is_constant); - ROCP_SDK_SAVE_DATA_BITFIELD("is_derived", is_derived); - ROCP_SDK_SAVE_DATA_CSTR(name); - ROCP_SDK_SAVE_DATA_CSTR(description); - ROCP_SDK_SAVE_DATA_CSTR(block); - ROCP_SDK_SAVE_DATA_CSTR(expression); - - auto convert = [](const auto* val, uint64_t sz) { - using data_type = std::remove_cv_t>; - auto retval = std::vector{}; - for(uint64_t i = 0; i < sz; ++i) - retval.emplace_back(val[i]); - return retval; - }; - - ROCP_SDK_SAVE_VALUE("dims", convert(data.dimensions, data.dimensions_count)); - ROCP_SDK_SAVE_VALUE("instances", convert(data.instance_ids, data.instance_ids_count)); -} - -template -void -save(ArchiveT& ar, rocprofiler_counter_record_dimension_info_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(id); - ROCP_SDK_SAVE_DATA_FIELD(instance_size); - ROCP_SDK_SAVE_DATA_CSTR(name); -} - -template -void -save(ArchiveT& ar, rocprofiler_callback_tracing_runtime_initialization_data_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(version); - ROCP_SDK_SAVE_DATA_FIELD(instance); -} - -template -void -save(ArchiveT& ar, rocprofiler_buffer_tracing_runtime_initialization_record_t data) -{ - ROCP_SDK_SAVE_DATA_FIELD(size); - ROCP_SDK_SAVE_DATA_FIELD(kind); - ROCP_SDK_SAVE_DATA_FIELD(operation); - ROCP_SDK_SAVE_DATA_FIELD(correlation_id); - ROCP_SDK_SAVE_DATA_FIELD(timestamp); - ROCP_SDK_SAVE_DATA_FIELD(thread_id); - ROCP_SDK_SAVE_DATA_FIELD(version); - ROCP_SDK_SAVE_DATA_FIELD(instance); -} - -template -void -save(ArchiveT& ar, const rocprofiler::sdk::utility::name_info& data) -{ - ar.makeArray(); - for(const auto& itr : data) - ar(cereal::make_nvp("entry", itr)); -} - -template -void -save(ArchiveT& ar, const rocprofiler::sdk::utility::name_info_impl& data) -{ - auto _name = std::string{data.name}; - auto _ops = std::vector{}; - _ops.reserve(data.operations.size()); - - ar(cereal::make_nvp("kind", _name)); - for(auto itr : data.operations) - _ops.emplace_back(itr); - ar(cereal::make_nvp("operations", _ops)); -} - -ROCPROFILER_SDK_CEREAL_NAMESPACE_END - -#undef ROCP_SDK_SAVE_DATA_FIELD -#undef ROCP_SDK_SAVE_DATA_VALUE -#undef ROCP_SDK_SAVE_DATA_CSTR -#undef ROCP_SDK_SAVE_DATA_BITFIELD +#include +#include diff --git a/source/include/rocprofiler-sdk/cxx/serialization/CMakeLists.txt b/source/include/rocprofiler-sdk/cxx/serialization/CMakeLists.txt new file mode 100644 index 0000000000..8701c73691 --- /dev/null +++ b/source/include/rocprofiler-sdk/cxx/serialization/CMakeLists.txt @@ -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) diff --git a/source/include/rocprofiler-sdk/cxx/serialization/load.hpp b/source/include/rocprofiler-sdk/cxx/serialization/load.hpp new file mode 100644 index 0000000000..da53dff5fb --- /dev/null +++ b/source/include/rocprofiler-sdk/cxx/serialization/load.hpp @@ -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 +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +#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 +void +load(ArchiveT& ar, rocprofiler_agent_id_t& data) +{ + ROCP_SDK_LOAD_DATA_FIELD(handle); +} + +template +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 +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 +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 +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 +void +load(ArchiveT& ar, HSA_MEMORYPROPERTY& data) +{ + ROCP_SDK_LOAD_DATA_BITFIELD("HotPluggable", ui32.HotPluggable); + ROCP_SDK_LOAD_DATA_BITFIELD("NonVolatile", ui32.NonVolatile); +} + +template +void +load(ArchiveT& ar, HSA_ENGINE_VERSION& data) +{ + ROCP_SDK_LOAD_DATA_BITFIELD("uCodeSDMA", uCodeSDMA); + ROCP_SDK_LOAD_DATA_BITFIELD("uCodeRes", uCodeRes); +} + +template +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 +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 +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 +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 +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>>; + auto vec = std::vector{}; + 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 +void +load(ArchiveT& ar, rocprofiler_address_t& data) +{ + ROCP_SDK_LOAD_DATA_FIELD(handle); +} + +template +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 +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 diff --git a/source/include/rocprofiler-sdk/cxx/serialization/save.hpp b/source/include/rocprofiler-sdk/cxx/serialization/save.hpp new file mode 100644 index 0000000000..97f2dae792 --- /dev/null +++ b/source/include/rocprofiler-sdk/cxx/serialization/save.hpp @@ -0,0 +1,1342 @@ +// 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 +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include + +#define ROCP_SDK_SAVE_DATA_FIELD(FIELD) ar(make_nvp(#FIELD, data.FIELD)) +#define ROCP_SDK_SAVE_DATA_VALUE(NAME, VALUE) ar(make_nvp(NAME, data.VALUE)) +#define ROCP_SDK_SAVE_VALUE(NAME, VALUE) ar(make_nvp(NAME, VALUE)) +#define ROCP_SDK_SAVE_DATA_CSTR(FIELD) \ + ar(make_nvp(#FIELD, std::string{data.FIELD ? data.FIELD : ""})) +#define ROCP_SDK_SAVE_DATA_BITFIELD(NAME, VALUE) \ + { \ + auto _val = data.VALUE; \ + ar(make_nvp(NAME, _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 + +namespace rocprofiler +{ +namespace sdk +{ +namespace serialization +{ +struct buffer_tracing_args +{ + std::string type = {}; + std::string name = {}; + std::string value = {}; +}; + +template +auto +get_buffer_tracing_args(Tp& data) +{ + auto populate_args_array = [](rocprofiler_buffer_tracing_kind_t /*kind*/, + rocprofiler_tracing_operation_t /*operation*/, + uint32_t arg_number, + const void* const /*arg_value_addr*/, + int32_t /*arg_indirection_count*/, + const char* arg_type, + const char* arg_name, + const char* arg_value_str, + void* cb_data) -> int { + if(!cb_data) return 1; + + auto* vec = static_cast*>(cb_data); + auto sz = std::max(arg_number + 1, vec->size()); + vec->resize(sz, buffer_tracing_args{}); + vec->at(arg_number) = buffer_tracing_args{arg_type, arg_name, arg_value_str}; + + return 0; + }; + + auto ret = std::vector{}; + auto record = rocprofiler_record_header_t{}; + record.hash = + rocprofiler_record_header_compute_hash(ROCPROFILER_BUFFER_CATEGORY_TRACING, data.kind); + record.payload = &data; + + rocprofiler_iterate_buffer_tracing_record_args(record, populate_args_array, &ret); + + return ret; +} +} // namespace serialization +} // namespace sdk +} // namespace rocprofiler + +ROCPROFILER_SDK_CEREAL_NAMESPACE_BEGIN + +namespace sdk = ::rocprofiler::sdk; + +template +void +save(ArchiveT& ar, const sdk::serialization::buffer_tracing_args& data) +{ + ROCP_SDK_SAVE_DATA_FIELD(type); + ROCP_SDK_SAVE_DATA_FIELD(name); + ROCP_SDK_SAVE_DATA_FIELD(value); +} + +template +void +save(ArchiveT& ar, rocprofiler_context_id_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(handle); +} + +template +void +save(ArchiveT& ar, rocprofiler_agent_id_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(handle); +} + +template +void +save(ArchiveT& ar, hsa_agent_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(handle); +} + +template +void +save(ArchiveT& ar, rocprofiler_queue_id_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(handle); +} + +template +void +save(ArchiveT& ar, rocprofiler_stream_id_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(handle); +} + +template +void +save(ArchiveT& ar, rocprofiler_counter_id_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(handle); +} + +template +void +save(ArchiveT& ar, rocprofiler_correlation_id_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(internal); + ROCP_SDK_SAVE_DATA_VALUE("external", external.value); + ROCP_SDK_SAVE_DATA_VALUE("ancestor", ancestor); +} + +template +void +save(ArchiveT& ar, rocprofiler_async_correlation_id_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(internal); + ROCP_SDK_SAVE_DATA_VALUE("external", external.value); +} + +template +void +save(ArchiveT& ar, rocprofiler_dim3_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(x); + ROCP_SDK_SAVE_DATA_FIELD(y); + ROCP_SDK_SAVE_DATA_FIELD(z); +} + +template +void +save(ArchiveT& ar, rocprofiler_address_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(handle); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_code_object_load_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(code_object_id); + ROCP_SDK_SAVE_DATA_FIELD(agent_id); + ROCP_SDK_SAVE_DATA_CSTR(uri); + ROCP_SDK_SAVE_DATA_FIELD(load_base); + ROCP_SDK_SAVE_DATA_FIELD(load_size); + ROCP_SDK_SAVE_DATA_FIELD(load_delta); + ROCP_SDK_SAVE_DATA_FIELD(storage_type); + if(data.storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_FILE) + { + ROCP_SDK_SAVE_DATA_FIELD(storage_file); + } + else if(data.storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_MEMORY) + { + ROCP_SDK_SAVE_DATA_FIELD(memory_base); + ROCP_SDK_SAVE_DATA_FIELD(memory_size); + } +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kernel_id); + ROCP_SDK_SAVE_DATA_FIELD(code_object_id); + ROCP_SDK_SAVE_DATA_CSTR(kernel_name); + ROCP_SDK_SAVE_DATA_FIELD(kernel_object); + ROCP_SDK_SAVE_DATA_FIELD(kernarg_segment_size); + ROCP_SDK_SAVE_DATA_FIELD(kernarg_segment_alignment); + ROCP_SDK_SAVE_DATA_FIELD(group_segment_size); + ROCP_SDK_SAVE_DATA_FIELD(private_segment_size); + ROCP_SDK_SAVE_DATA_FIELD(sgpr_count); + ROCP_SDK_SAVE_DATA_FIELD(arch_vgpr_count); + ROCP_SDK_SAVE_DATA_FIELD(accum_vgpr_count); + ROCP_SDK_SAVE_DATA_FIELD(kernel_code_entry_byte_offset); + ROCP_SDK_SAVE_DATA_FIELD(kernel_address); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_code_object_host_kernel_symbol_register_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(host_function_id); + ROCP_SDK_SAVE_DATA_FIELD(kernel_id); + ROCP_SDK_SAVE_DATA_FIELD(code_object_id); + ROCP_SDK_SAVE_DATA_FIELD(host_function); + ROCP_SDK_SAVE_DATA_FIELD(modules); + ROCP_SDK_SAVE_DATA_CSTR(device_function); + ROCP_SDK_SAVE_DATA_FIELD(thread_limit); + ROCP_SDK_SAVE_DATA_FIELD(thread_ids); + ROCP_SDK_SAVE_DATA_FIELD(block_ids); + ROCP_SDK_SAVE_DATA_FIELD(block_dims); + ROCP_SDK_SAVE_DATA_FIELD(grid_dims); + ROCP_SDK_SAVE_DATA_FIELD(workgroup_size); +} + +template +void +save(ArchiveT& ar, rocprofiler_hsa_api_retval_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(uint64_t_retval); +} + +template +void +save(ArchiveT& ar, const hsa_queue_t& data) +{ + ar(make_nvp("queue_id", data.id)); +} + +template +void +save(ArchiveT& ar, hsa_amd_event_scratch_alloc_start_t data) +{ + ar(make_nvp("queue_id", *data.queue)); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); +} + +template +void +save(ArchiveT& ar, hsa_amd_event_scratch_alloc_end_t data) +{ + ar(make_nvp("queue_id", *data.queue)); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(num_slots); + ROCP_SDK_SAVE_DATA_FIELD(flags); +} + +template +void +save(ArchiveT& ar, hsa_amd_event_scratch_free_start_t data) +{ + ar(make_nvp("queue_id", *data.queue)); +} + +template +void +save(ArchiveT& ar, hsa_amd_event_scratch_free_end_t data) +{ + ar(make_nvp("queue_id", *data.queue)); + ROCP_SDK_SAVE_DATA_FIELD(flags); +} + +template +void +save(ArchiveT& ar, hsa_amd_event_scratch_async_reclaim_start_t data) +{ + ar(make_nvp("queue_id", *data.queue)); +} + +template +void +save(ArchiveT& ar, hsa_amd_event_scratch_async_reclaim_end_t data) +{ + ar(make_nvp("queue_id", *data.queue)); + ROCP_SDK_SAVE_DATA_FIELD(flags); +} + +template +void +save(ArchiveT& ar, rocprofiler_marker_api_retval_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(int64_t_retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_hsa_api_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + // ROCP_SDK_SAVE_DATA_FIELD(args); + ROCP_SDK_SAVE_DATA_FIELD(retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_marker_api_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + // ROCP_SDK_SAVE_DATA_FIELD(args); + ROCP_SDK_SAVE_DATA_FIELD(retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_hip_api_retval_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(uint64_t_retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_hip_api_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + // ROCP_SDK_SAVE_DATA_FIELD(args); + ROCP_SDK_SAVE_DATA_FIELD(retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_scratch_memory_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(agent_id); + ROCP_SDK_SAVE_DATA_FIELD(queue_id); + ROCP_SDK_SAVE_DATA_FIELD(flags); + ROCP_SDK_SAVE_DATA_FIELD(args_kind); +} + +template +void +save(ArchiveT& ar, rocprofiler_kernel_dispatch_info_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(agent_id); + ROCP_SDK_SAVE_DATA_FIELD(queue_id); + ROCP_SDK_SAVE_DATA_FIELD(kernel_id); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); + ROCP_SDK_SAVE_DATA_FIELD(private_segment_size); + ROCP_SDK_SAVE_DATA_FIELD(group_segment_size); + ROCP_SDK_SAVE_DATA_FIELD(workgroup_size); + ROCP_SDK_SAVE_DATA_FIELD(grid_size); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_kernel_dispatch_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_info); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_memory_copy_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(dst_agent_id); + ROCP_SDK_SAVE_DATA_FIELD(src_agent_id); + ROCP_SDK_SAVE_DATA_FIELD(bytes); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_memory_allocation_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(agent_id); + ROCP_SDK_SAVE_VALUE("address", rocprofiler::sdk::utility::as_hex(data.address.handle, 16)); + ROCP_SDK_SAVE_DATA_FIELD(allocation_size); +} + +template +void +save(ArchiveT& ar, rocprofiler_rccl_api_retval_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(ncclResult_t_retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_rccl_api_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + // ROCP_SDK_SAVE_DATA_FIELD(args); + ROCP_SDK_SAVE_DATA_FIELD(retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_rocdecode_api_retval_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(uint64_t_retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_rocdecode_api_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_rocjpeg_api_retval_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(rocJpegStatus_retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_rocjpeg_api_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_ompt_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + // ROCP_SDK_SAVE_DATA_FIELD(args); +} + +template +void +save(ArchiveT& ar, rocprofiler_dispatch_counting_service_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_info); +} + +template +void +save(ArchiveT& ar, rocprofiler_dispatch_counting_service_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(num_records); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_info); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(context_id); + ROCP_SDK_SAVE_DATA_FIELD(thread_id); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(operation); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(phase); +} + +template +void +save_buffer_tracing_api_record(ArchiveT& ar, Tp data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(operation); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(thread_id); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_hsa_api_record_t data) +{ + save_buffer_tracing_api_record(ar, data); +} + +template +void +save(ArchiveT& ar, rocprofiler_counter_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(id); + ROCP_SDK_SAVE_DATA_FIELD(counter_value); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_hip_api_record_t data) +{ + save_buffer_tracing_api_record(ar, data); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_hip_api_ext_record_t data) +{ + save_buffer_tracing_api_record(ar, data); + auto args = sdk::serialization::get_buffer_tracing_args(data); + ROCP_SDK_SAVE_VALUE("args", args); + ROCP_SDK_SAVE_DATA_FIELD(retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_marker_api_record_t data) +{ + save_buffer_tracing_api_record(ar, data); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_rccl_api_record_t data) +{ + save_buffer_tracing_api_record(ar, data); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_rocdecode_api_record_t data) +{ + save_buffer_tracing_api_record(ar, data); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_rocdecode_api_ext_record_t data) +{ + save_buffer_tracing_api_record(ar, data); + auto args = sdk::serialization::get_buffer_tracing_args(data); + ROCP_SDK_SAVE_VALUE("args", args); + ROCP_SDK_SAVE_DATA_FIELD(retval); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_rocjpeg_api_record_t data) +{ + save_buffer_tracing_api_record(ar, data); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_ompt_target_t data) +{ + ROCP_SDK_SAVE_DATA_VALUE("kind", kind); + ROCP_SDK_SAVE_DATA_VALUE("device", device_num); + ROCP_SDK_SAVE_DATA_VALUE("task_id", task_id); + ROCP_SDK_SAVE_DATA_VALUE("target_id", target_id); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_ompt_target_data_op_t data) +{ + ROCP_SDK_SAVE_DATA_VALUE("host_op_id", host_op_id); + ROCP_SDK_SAVE_DATA_VALUE("optype", optype); + ROCP_SDK_SAVE_DATA_VALUE("src_device_num", src_device_num); + ROCP_SDK_SAVE_DATA_VALUE("dst_device_num", dst_device_num); + ROCP_SDK_SAVE_DATA_VALUE("bytes", bytes); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_ompt_target_kernel_t data) +{ + ROCP_SDK_SAVE_DATA_VALUE("host_op_id", host_op_id); + ROCP_SDK_SAVE_DATA_VALUE("device_num", device_num); + ROCP_SDK_SAVE_DATA_VALUE("requested_num_teams", requested_num_teams); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_ompt_record_t data) +{ + save_buffer_tracing_api_record(ar, data); + + if(data.operation == ROCPROFILER_OMPT_ID_target_emi) + { + ROCP_SDK_SAVE_DATA_FIELD(target); + } + else if(data.operation == ROCPROFILER_OMPT_ID_target_data_op_emi) + { + ROCP_SDK_SAVE_DATA_FIELD(target_data_op); + } + else if(data.operation == ROCPROFILER_OMPT_ID_target_submit_emi) + { + ROCP_SDK_SAVE_DATA_FIELD(target_kernel); + } +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_kernel_dispatch_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(operation); + ROCP_SDK_SAVE_DATA_FIELD(thread_id); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_info); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_memory_copy_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(operation); + ROCP_SDK_SAVE_DATA_FIELD(thread_id); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(dst_agent_id); + ROCP_SDK_SAVE_DATA_FIELD(src_agent_id); + ROCP_SDK_SAVE_DATA_FIELD(bytes); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_memory_allocation_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(operation); + ROCP_SDK_SAVE_DATA_FIELD(thread_id); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(agent_id); + ROCP_SDK_SAVE_VALUE("address", rocprofiler::sdk::utility::as_hex(data.address.handle, 16)); + ROCP_SDK_SAVE_DATA_FIELD(allocation_size); +} + +template +void +save(ArchiveT& ar, const rocprofiler_page_migration_page_fault_start_t& data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("read_fault", read_fault); + ROCP_SDK_SAVE_DATA_FIELD(agent_id); + ROCP_SDK_SAVE_DATA_FIELD(address); +} + +template +void +save(ArchiveT& ar, const rocprofiler_page_migration_page_fault_end_t& data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("migrated", migrated); + ROCP_SDK_SAVE_DATA_FIELD(agent_id); + ROCP_SDK_SAVE_DATA_FIELD(address); +} + +template +void +save(ArchiveT& ar, const rocprofiler_page_migration_page_migrate_start_t& data) +{ + ROCP_SDK_SAVE_DATA_FIELD(start_addr); + ROCP_SDK_SAVE_DATA_FIELD(end_addr); + ROCP_SDK_SAVE_DATA_FIELD(from_agent); + ROCP_SDK_SAVE_DATA_FIELD(to_agent); + ROCP_SDK_SAVE_DATA_FIELD(prefetch_agent); + ROCP_SDK_SAVE_DATA_FIELD(preferred_agent); + ROCP_SDK_SAVE_DATA_FIELD(trigger); +} + +template +void +save(ArchiveT& ar, const rocprofiler_page_migration_page_migrate_end_t& data) +{ + ROCP_SDK_SAVE_DATA_FIELD(start_addr); + ROCP_SDK_SAVE_DATA_FIELD(end_addr); + ROCP_SDK_SAVE_DATA_FIELD(from_agent); + ROCP_SDK_SAVE_DATA_FIELD(to_agent); + ROCP_SDK_SAVE_DATA_FIELD(trigger); + ROCP_SDK_SAVE_DATA_FIELD(error_code); +} + +template +void +save(ArchiveT& ar, const rocprofiler_page_migration_queue_eviction_t& data) +{ + ROCP_SDK_SAVE_DATA_FIELD(agent_id); + ROCP_SDK_SAVE_DATA_FIELD(trigger); +} + +template +void +save(ArchiveT& ar, const rocprofiler_page_migration_queue_restore_t& data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("rescheduled", rescheduled); + ROCP_SDK_SAVE_DATA_FIELD(agent_id); +} + +template +void +save(ArchiveT& ar, const rocprofiler_page_migration_unmap_from_gpu_t& data) +{ + ROCP_SDK_SAVE_DATA_FIELD(start_addr); + ROCP_SDK_SAVE_DATA_FIELD(end_addr); + ROCP_SDK_SAVE_DATA_FIELD(agent_id); + ROCP_SDK_SAVE_DATA_FIELD(trigger); +} + +template +void +save(ArchiveT& ar, const rocprofiler_page_migration_dropped_event_t& data) +{ + ROCP_SDK_SAVE_DATA_FIELD(dropped_events_count); +} + +namespace details +{ +template +struct save_page_migration_arg; + +#define ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(ENUM_VALUE, UNION_ARG) \ + template <> \ + struct save_page_migration_arg \ + { \ + static constexpr auto value = ROCPROFILER_PAGE_MIGRATION_##ENUM_VALUE; \ + template \ + void operator()(ArchiveT& ar, rocprofiler_page_migration_args_t args) \ + { \ + ar(make_nvp(#UNION_ARG, args.UNION_ARG)); \ + } \ + }; + +ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(NONE, none) +ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(PAGE_MIGRATE_START, page_migrate_start) +ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(PAGE_MIGRATE_END, page_migrate_end) +ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(PAGE_FAULT_START, page_fault_start) +ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(PAGE_FAULT_END, page_fault_end) +ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(QUEUE_EVICTION, queue_eviction) +ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(QUEUE_RESTORE, queue_restore) +ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(UNMAP_FROM_GPU, unmap_from_gpu) +ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG(DROPPED_EVENT, dropped_event) + +#undef ROCP_SDK_SPECIALIZE_PAGE_MIGRATION_ARG + +template +void +save_page_migration_args(ArchiveT& ar, + rocprofiler_page_migration_operation_t op, + rocprofiler_page_migration_args_t args, + std::index_sequence) +{ + using save_page_migration_type = save_page_migration_arg; + if(op == save_page_migration_type::value) + { + if constexpr(save_page_migration_type::value != ROCPROFILER_PAGE_MIGRATION_NONE) + save_page_migration_type{}(ar, args); + } + else if constexpr(sizeof...(IdxTail) > 0) + { + save_page_migration_args(ar, op, args, std::index_sequence{}); + } +} +} // namespace details + +template +void +save(ArchiveT& ar, const rocprofiler_buffer_tracing_page_migration_record_t& data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(operation); + ROCP_SDK_SAVE_DATA_FIELD(timestamp); + ROCP_SDK_SAVE_DATA_FIELD(pid); + details::save_page_migration_args( + ar, data.operation, data.args, std::make_index_sequence{}); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_scratch_memory_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(operation); + ROCP_SDK_SAVE_DATA_FIELD(agent_id); + ROCP_SDK_SAVE_DATA_FIELD(queue_id); + ROCP_SDK_SAVE_DATA_FIELD(thread_id); + ROCP_SDK_SAVE_DATA_FIELD(start_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(end_timestamp); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(flags); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_correlation_id_retirement_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(timestamp); + ROCP_SDK_SAVE_DATA_FIELD(internal_correlation_id); +} + +template +void +save(ArchiveT& ar, HsaCacheType data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("Data", ui32.Data); + ROCP_SDK_SAVE_DATA_BITFIELD("Instruction", ui32.Instruction); + ROCP_SDK_SAVE_DATA_BITFIELD("CPU", ui32.CPU); + ROCP_SDK_SAVE_DATA_BITFIELD("HSACU", ui32.HSACU); +} + +template +void +save(ArchiveT& ar, HSA_LINKPROPERTY data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("Override", ui32.Override); + ROCP_SDK_SAVE_DATA_BITFIELD("NonCoherent", ui32.NonCoherent); + ROCP_SDK_SAVE_DATA_BITFIELD("NoAtomics32bit", ui32.NoAtomics32bit); + ROCP_SDK_SAVE_DATA_BITFIELD("NoAtomics64bit", ui32.NoAtomics64bit); + ROCP_SDK_SAVE_DATA_BITFIELD("NoPeerToPeerDMA", ui32.NoPeerToPeerDMA); +} + +template +void +save(ArchiveT& ar, HSA_CAPABILITY data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("HotPluggable", ui32.HotPluggable); + ROCP_SDK_SAVE_DATA_BITFIELD("HSAMMUPresent", ui32.HSAMMUPresent); + ROCP_SDK_SAVE_DATA_BITFIELD("SharedWithGraphics", ui32.SharedWithGraphics); + ROCP_SDK_SAVE_DATA_BITFIELD("QueueSizePowerOfTwo", ui32.QueueSizePowerOfTwo); + ROCP_SDK_SAVE_DATA_BITFIELD("QueueSize32bit", ui32.QueueSize32bit); + ROCP_SDK_SAVE_DATA_BITFIELD("QueueIdleEvent", ui32.QueueIdleEvent); + ROCP_SDK_SAVE_DATA_BITFIELD("VALimit", ui32.VALimit); + ROCP_SDK_SAVE_DATA_BITFIELD("WatchPointsSupported", ui32.WatchPointsSupported); + ROCP_SDK_SAVE_DATA_BITFIELD("WatchPointsTotalBits", ui32.WatchPointsTotalBits); + ROCP_SDK_SAVE_DATA_BITFIELD("DoorbellType", ui32.DoorbellType); + ROCP_SDK_SAVE_DATA_BITFIELD("AQLQueueDoubleMap", ui32.AQLQueueDoubleMap); + ROCP_SDK_SAVE_DATA_BITFIELD("DebugTrapSupported", ui32.DebugTrapSupported); + ROCP_SDK_SAVE_DATA_BITFIELD("WaveLaunchTrapOverrideSupported", + ui32.WaveLaunchTrapOverrideSupported); + ROCP_SDK_SAVE_DATA_BITFIELD("WaveLaunchModeSupported", ui32.WaveLaunchModeSupported); + ROCP_SDK_SAVE_DATA_BITFIELD("PreciseMemoryOperationsSupported", + ui32.PreciseMemoryOperationsSupported); + ROCP_SDK_SAVE_DATA_BITFIELD("DEPRECATED_SRAM_EDCSupport", ui32.DEPRECATED_SRAM_EDCSupport); + ROCP_SDK_SAVE_DATA_BITFIELD("Mem_EDCSupport", ui32.Mem_EDCSupport); + ROCP_SDK_SAVE_DATA_BITFIELD("RASEventNotify", ui32.RASEventNotify); + ROCP_SDK_SAVE_DATA_BITFIELD("ASICRevision", ui32.ASICRevision); + ROCP_SDK_SAVE_DATA_BITFIELD("SRAM_EDCSupport", ui32.SRAM_EDCSupport); + ROCP_SDK_SAVE_DATA_BITFIELD("SVMAPISupported", ui32.SVMAPISupported); + ROCP_SDK_SAVE_DATA_BITFIELD("CoherentHostAccess", ui32.CoherentHostAccess); + ROCP_SDK_SAVE_DATA_BITFIELD("DebugSupportedFirmware", ui32.DebugSupportedFirmware); +} + +template +void +save(ArchiveT& ar, HSA_MEMORYPROPERTY data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("HotPluggable", ui32.HotPluggable); + ROCP_SDK_SAVE_DATA_BITFIELD("NonVolatile", ui32.NonVolatile); +} + +template +void +save(ArchiveT& ar, HSA_ENGINE_VERSION data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("uCodeSDMA", uCodeSDMA); + ROCP_SDK_SAVE_DATA_BITFIELD("uCodeRes", uCodeRes); +} + +template +void +save(ArchiveT& ar, HSA_ENGINE_ID data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("uCode", ui32.uCode); + ROCP_SDK_SAVE_DATA_BITFIELD("Major", ui32.Major); + ROCP_SDK_SAVE_DATA_BITFIELD("Minor", ui32.Minor); + ROCP_SDK_SAVE_DATA_BITFIELD("Stepping", ui32.Stepping); +} + +template +void +save(ArchiveT& ar, rocprofiler_agent_cache_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(processor_id_low); + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(level); + ROCP_SDK_SAVE_DATA_FIELD(cache_line_size); + ROCP_SDK_SAVE_DATA_FIELD(cache_lines_per_tag); + ROCP_SDK_SAVE_DATA_FIELD(association); + ROCP_SDK_SAVE_DATA_FIELD(latency); + ROCP_SDK_SAVE_DATA_FIELD(type); +} +template +void +save(ArchiveT& ar, rocprofiler_pc_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(code_object_id); + ROCP_SDK_SAVE_DATA_FIELD(code_object_offset); +} + +template +void +save(ArchiveT& ar, rocprofiler_pc_sampling_hw_id_v0_t data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("chiplet", chiplet); + ROCP_SDK_SAVE_DATA_BITFIELD("wave_id", wave_id); + ROCP_SDK_SAVE_DATA_BITFIELD("simd_id", simd_id); + ROCP_SDK_SAVE_DATA_BITFIELD("pipe_id", pipe_id); + ROCP_SDK_SAVE_DATA_BITFIELD("cu_or_wgp_id", cu_or_wgp_id); + ROCP_SDK_SAVE_DATA_BITFIELD("shader_array_id", shader_array_id); + ROCP_SDK_SAVE_DATA_BITFIELD("shader_engine_id", shader_engine_id); + ROCP_SDK_SAVE_DATA_BITFIELD("workgroup_id ", workgroup_id); + ROCP_SDK_SAVE_DATA_BITFIELD("vm_id", vm_id); + ROCP_SDK_SAVE_DATA_BITFIELD("queue_id", queue_id); + ROCP_SDK_SAVE_DATA_BITFIELD("microengine_id", microengine_id); +} + +template +void +save(ArchiveT& ar, rocprofiler_pc_sampling_record_host_trap_v0_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(hw_id); + ROCP_SDK_SAVE_DATA_FIELD(pc); + ROCP_SDK_SAVE_DATA_FIELD(exec_mask); + ROCP_SDK_SAVE_DATA_FIELD(timestamp); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); + ROCP_SDK_SAVE_DATA_VALUE("corr_id", correlation_id); + ROCP_SDK_SAVE_DATA_VALUE("wrkgrp_id", workgroup_id); + ROCP_SDK_SAVE_DATA_BITFIELD("wave_in_grp", wave_in_group); +} + +template +void +save(ArchiveT& ar, rocprofiler_pc_sampling_record_stochastic_header_t data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("has_mem_cnt", has_memory_counter); +} + +template +void +save_pc_sampling_inst_type(ArchiveT& ar, rocprofiler_pc_sampling_instruction_type_t inst_type) +{ + ar(make_nvp("inst_type", + std::string(rocprofiler_get_pc_sampling_instruction_type_name(inst_type)))); +} + +template +void +save_pc_sampling_stall_reason(ArchiveT& ar, + rocprofiler_pc_sampling_instruction_not_issued_reason_t stall_reason) +{ + ar(make_nvp( + "stall_reason", + std::string(rocprofiler_get_pc_sampling_instruction_not_issued_reason_name(stall_reason)))); +} + +template +void +save(ArchiveT& ar, rocprofiler_pc_sampling_snapshot_v0_t data) +{ + save_pc_sampling_stall_reason( + ar, + static_cast( + data.reason_not_issued)); + + ROCP_SDK_SAVE_DATA_BITFIELD("dual_issue_valu", dual_issue_valu); + + // Arb state (pipe issued) + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_valu", arb_state_issue_valu); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_matrix", arb_state_issue_matrix); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_lds", arb_state_issue_lds); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_lds_direct", arb_state_issue_lds_direct); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_scalar", arb_state_issue_scalar); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_vmem_tex", arb_state_issue_vmem_tex); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_flat", arb_state_issue_flat); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_exp", arb_state_issue_exp); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_misc", arb_state_issue_misc); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_issue_brmsg", arb_state_issue_brmsg); + // Arb state (pipe stalled) + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_valu", arb_state_stall_valu); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_matrix", arb_state_stall_matrix); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_lds", arb_state_stall_lds); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_lds_direct", arb_state_stall_lds_direct); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_scalar", arb_state_stall_scalar); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_vmem_tex", arb_state_stall_vmem_tex); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_flat", arb_state_stall_flat); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_exp", arb_state_stall_exp); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_misc", arb_state_stall_misc); + ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_brmsg", arb_state_stall_brmsg); +} + +template +void +save(ArchiveT& ar, rocprofiler_pc_sampling_record_stochastic_v0_t data) +{ + // flags specific for stochastic sampling + ROCP_SDK_SAVE_DATA_FIELD(flags); + + // Common for host-trap and stochastic + ROCP_SDK_SAVE_DATA_FIELD(hw_id); + ROCP_SDK_SAVE_DATA_FIELD(pc); + ROCP_SDK_SAVE_DATA_FIELD(exec_mask); + ROCP_SDK_SAVE_DATA_FIELD(timestamp); + ROCP_SDK_SAVE_DATA_FIELD(dispatch_id); + ROCP_SDK_SAVE_DATA_VALUE("corr_id", correlation_id); + ROCP_SDK_SAVE_DATA_VALUE("wrkgrp_id", workgroup_id); + ROCP_SDK_SAVE_DATA_BITFIELD("wave_in_grp", wave_in_group); + + // fields specific for stochastic + ROCP_SDK_SAVE_DATA_BITFIELD("wave_issued", wave_issued); + save_pc_sampling_inst_type( + ar, static_cast(data.inst_type)); + ROCP_SDK_SAVE_DATA_BITFIELD("wave_cnt", wave_count); + ROCP_SDK_SAVE_DATA_FIELD(snapshot); + + // TODO: add memory counters +} + +template +void +save(ArchiveT& ar, rocprofiler_agent_io_link_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(type); + ROCP_SDK_SAVE_DATA_FIELD(version_major); + ROCP_SDK_SAVE_DATA_FIELD(version_minor); + ROCP_SDK_SAVE_DATA_FIELD(node_from); + ROCP_SDK_SAVE_DATA_FIELD(node_to); + ROCP_SDK_SAVE_DATA_FIELD(weight); + ROCP_SDK_SAVE_DATA_FIELD(min_latency); + ROCP_SDK_SAVE_DATA_FIELD(max_latency); + ROCP_SDK_SAVE_DATA_FIELD(min_bandwidth); + ROCP_SDK_SAVE_DATA_FIELD(max_bandwidth); + ROCP_SDK_SAVE_DATA_FIELD(recommended_transfer_size); + ROCP_SDK_SAVE_DATA_FIELD(flags); +} + +template +void +save(ArchiveT& ar, rocprofiler_agent_mem_bank_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(heap_type); + ROCP_SDK_SAVE_DATA_FIELD(flags); + ROCP_SDK_SAVE_DATA_FIELD(width); + ROCP_SDK_SAVE_DATA_FIELD(mem_clk_max); + ROCP_SDK_SAVE_DATA_FIELD(size_in_bytes); +} + +template +void +save(ArchiveT& ar, rocprofiler_pc_sampling_configuration_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(method); + ROCP_SDK_SAVE_DATA_FIELD(unit); + ROCP_SDK_SAVE_DATA_FIELD(min_interval); + ROCP_SDK_SAVE_DATA_FIELD(max_interval); + ROCP_SDK_SAVE_DATA_FIELD(flags); +} + +template +void +save(ArchiveT& ar, rocprofiler_agent_runtime_visiblity_t data) +{ + ROCP_SDK_SAVE_DATA_BITFIELD("hsa", hsa); + ROCP_SDK_SAVE_DATA_BITFIELD("hip", hip); + ROCP_SDK_SAVE_DATA_BITFIELD("rccl", rccl); + ROCP_SDK_SAVE_DATA_BITFIELD("rocdecode", rocdecode); +} + +template +void +save(ArchiveT& ar, rocprofiler_uuid_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(value); +} + +template +void +save(ArchiveT& ar, const rocprofiler_agent_v0_t& data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(id); + ROCP_SDK_SAVE_DATA_FIELD(type); + ROCP_SDK_SAVE_DATA_FIELD(cpu_cores_count); + ROCP_SDK_SAVE_DATA_FIELD(simd_count); + ROCP_SDK_SAVE_DATA_FIELD(mem_banks_count); + ROCP_SDK_SAVE_DATA_FIELD(caches_count); + ROCP_SDK_SAVE_DATA_FIELD(io_links_count); + ROCP_SDK_SAVE_DATA_FIELD(cpu_core_id_base); + ROCP_SDK_SAVE_DATA_FIELD(simd_id_base); + ROCP_SDK_SAVE_DATA_FIELD(max_waves_per_simd); + ROCP_SDK_SAVE_DATA_FIELD(lds_size_in_kb); + ROCP_SDK_SAVE_DATA_FIELD(gds_size_in_kb); + ROCP_SDK_SAVE_DATA_FIELD(num_gws); + ROCP_SDK_SAVE_DATA_FIELD(wave_front_size); + ROCP_SDK_SAVE_DATA_FIELD(num_xcc); + ROCP_SDK_SAVE_DATA_FIELD(cu_count); + ROCP_SDK_SAVE_DATA_FIELD(array_count); + ROCP_SDK_SAVE_DATA_FIELD(num_shader_banks); + ROCP_SDK_SAVE_DATA_FIELD(simd_arrays_per_engine); + ROCP_SDK_SAVE_DATA_FIELD(cu_per_simd_array); + ROCP_SDK_SAVE_DATA_FIELD(simd_per_cu); + ROCP_SDK_SAVE_DATA_FIELD(max_slots_scratch_cu); + ROCP_SDK_SAVE_DATA_FIELD(gfx_target_version); + ROCP_SDK_SAVE_DATA_FIELD(vendor_id); + ROCP_SDK_SAVE_DATA_FIELD(device_id); + ROCP_SDK_SAVE_DATA_FIELD(location_id); + ROCP_SDK_SAVE_DATA_FIELD(domain); + ROCP_SDK_SAVE_DATA_FIELD(drm_render_minor); + ROCP_SDK_SAVE_DATA_FIELD(num_sdma_engines); + ROCP_SDK_SAVE_DATA_FIELD(num_sdma_xgmi_engines); + ROCP_SDK_SAVE_DATA_FIELD(num_sdma_queues_per_engine); + ROCP_SDK_SAVE_DATA_FIELD(num_cp_queues); + ROCP_SDK_SAVE_DATA_FIELD(max_engine_clk_ccompute); + ROCP_SDK_SAVE_DATA_FIELD(max_engine_clk_fcompute); + ROCP_SDK_SAVE_DATA_FIELD(sdma_fw_version); + ROCP_SDK_SAVE_DATA_FIELD(fw_version); + ROCP_SDK_SAVE_DATA_FIELD(capability); + ROCP_SDK_SAVE_DATA_FIELD(cu_per_engine); + ROCP_SDK_SAVE_DATA_FIELD(max_waves_per_cu); + ROCP_SDK_SAVE_DATA_FIELD(family_id); + ROCP_SDK_SAVE_DATA_FIELD(workgroup_max_size); + ROCP_SDK_SAVE_DATA_FIELD(grid_max_size); + ROCP_SDK_SAVE_DATA_FIELD(local_mem_size); + ROCP_SDK_SAVE_DATA_FIELD(hive_id); + ROCP_SDK_SAVE_DATA_FIELD(gpu_id); + ROCP_SDK_SAVE_DATA_FIELD(workgroup_max_dim); + ROCP_SDK_SAVE_DATA_FIELD(grid_max_dim); + ROCP_SDK_SAVE_DATA_CSTR(name); + ROCP_SDK_SAVE_DATA_CSTR(vendor_name); + ROCP_SDK_SAVE_DATA_CSTR(product_name); + ROCP_SDK_SAVE_DATA_CSTR(model_name); + ROCP_SDK_SAVE_DATA_FIELD(node_id); + ROCP_SDK_SAVE_DATA_FIELD(logical_node_id); + ROCP_SDK_SAVE_DATA_FIELD(logical_node_type_id); + ROCP_SDK_SAVE_DATA_FIELD(runtime_visibility); + ROCP_SDK_SAVE_DATA_FIELD(uuid); + + auto generate = [&](auto name, const auto* value, uint64_t size) { + using value_type = std::remove_const_t>; + auto vec = std::vector{}; + vec.reserve(size); + for(uint64_t i = 0; i < size; ++i) + vec.emplace_back(value[i]); + ar(make_nvp(name, vec)); + }; + + 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 +void +save(ArchiveT& ar, rocprofiler_counter_info_v0_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(id); + ROCP_SDK_SAVE_DATA_BITFIELD("is_constant", is_constant); + ROCP_SDK_SAVE_DATA_BITFIELD("is_derived", is_derived); + ROCP_SDK_SAVE_DATA_CSTR(name); + ROCP_SDK_SAVE_DATA_CSTR(description); + ROCP_SDK_SAVE_DATA_CSTR(block); + ROCP_SDK_SAVE_DATA_CSTR(expression); +} + +template +void +save(ArchiveT& ar, rocprofiler_counter_info_v1_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(id); + ROCP_SDK_SAVE_DATA_BITFIELD("is_constant", is_constant); + ROCP_SDK_SAVE_DATA_BITFIELD("is_derived", is_derived); + ROCP_SDK_SAVE_DATA_CSTR(name); + ROCP_SDK_SAVE_DATA_CSTR(description); + ROCP_SDK_SAVE_DATA_CSTR(block); + ROCP_SDK_SAVE_DATA_CSTR(expression); + + auto convert = [](const auto* val, uint64_t sz) { + using data_type = std::remove_cv_t>; + auto retval = std::vector{}; + for(uint64_t i = 0; i < sz; ++i) + retval.emplace_back(val[i]); + return retval; + }; + + ROCP_SDK_SAVE_VALUE("dims", convert(data.dimensions, data.dimensions_count)); + ROCP_SDK_SAVE_VALUE("instances", convert(data.instance_ids, data.instance_ids_count)); +} + +template +void +save(ArchiveT& ar, rocprofiler_counter_record_dimension_info_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(id); + ROCP_SDK_SAVE_DATA_FIELD(instance_size); + ROCP_SDK_SAVE_DATA_CSTR(name); +} + +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_runtime_initialization_data_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(version); + ROCP_SDK_SAVE_DATA_FIELD(instance); +} + +template +void +save(ArchiveT& ar, rocprofiler_buffer_tracing_runtime_initialization_record_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(size); + ROCP_SDK_SAVE_DATA_FIELD(kind); + ROCP_SDK_SAVE_DATA_FIELD(operation); + ROCP_SDK_SAVE_DATA_FIELD(correlation_id); + ROCP_SDK_SAVE_DATA_FIELD(timestamp); + ROCP_SDK_SAVE_DATA_FIELD(thread_id); + ROCP_SDK_SAVE_DATA_FIELD(version); + ROCP_SDK_SAVE_DATA_FIELD(instance); +} + +template +void +save(ArchiveT& ar, const rocprofiler::sdk::utility::name_info& data) +{ + ar.makeArray(); + for(const auto& itr : data) + ar(cereal::make_nvp("entry", itr)); +} + +template +void +save(ArchiveT& ar, const rocprofiler::sdk::utility::name_info_impl& data) +{ + auto _name = std::string{data.name}; + auto _ops = std::vector{}; + _ops.reserve(data.operations.size()); + + ar(cereal::make_nvp("kind", _name)); + for(auto itr : data.operations) + _ops.emplace_back(itr); + ar(cereal::make_nvp("operations", _ops)); +} + +ROCPROFILER_SDK_CEREAL_NAMESPACE_END + +#undef ROCP_SDK_SAVE_DATA_FIELD +#undef ROCP_SDK_SAVE_DATA_VALUE +#undef ROCP_SDK_SAVE_DATA_CSTR +#undef ROCP_SDK_SAVE_DATA_BITFIELD diff --git a/source/include/rocprofiler-sdk/intercept_table.h b/source/include/rocprofiler-sdk/intercept_table.h index 06f5292f21..f7813dca09 100644 --- a/source/include/rocprofiler-sdk/intercept_table.h +++ b/source/include/rocprofiler-sdk/intercept_table.h @@ -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 diff --git a/source/include/rocprofiler-sdk/ompt/omp-tools.h b/source/include/rocprofiler-sdk/ompt/omp-tools.h index 1ac41d6136..74b9fec2d8 100644 --- a/source/include/rocprofiler-sdk/ompt/omp-tools.h +++ b/source/include/rocprofiler-sdk/ompt/omp-tools.h @@ -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 */ \ diff --git a/source/lib/common/CMakeLists.txt b/source/lib/common/CMakeLists.txt index 81e1c83773..4fd1182885 100644 --- a/source/lib/common/CMakeLists.txt +++ b/source/lib/common/CMakeLists.txt @@ -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 diff --git a/source/lib/common/hasher.hpp b/source/lib/common/hasher.hpp new file mode 100644 index 0000000000..22b0b2326d --- /dev/null +++ b/source/lib/common/hasher.hpp @@ -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 + +#include +#include +#include +#include + +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 ::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 + 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 + 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 ::value, bool>> +void +fnv1a_hasher::update(Tp data) +{ + update(reinterpret_cast(&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(data[i]); + // Note: Arithmetic overflow of unsigned integers is well defined in C++ standard unlike + // signed integers. + m_result *= kFnv1a64Prime; + } +} + +template +uint64_t +fnv1a_hasher::combine(Tp&& arg, TailT&&... args) +{ + auto _hasher = fnv1a_hasher{}; + _hasher.update_all(std::forward(arg), std::forward(args)...); + return _hasher.digest(); +} + +template +void +fnv1a_hasher::update_all(Tp&& arg, TailT&&... args) +{ + update(arg); + if constexpr(sizeof...(TailT) > 0) update_all(std::forward(args)...); +} +} // namespace common +} // namespace rocprofiler diff --git a/source/lib/common/md5sum.cpp b/source/lib/common/md5sum.cpp new file mode 100644 index 0000000000..6ab05369c5 --- /dev/null +++ b/source/lib/common/md5sum.cpp @@ -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 +#include +#include + +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(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 diff --git a/source/lib/common/md5sum.hpp b/source/lib/common/md5sum.hpp new file mode 100644 index 0000000000..49b384d8fa --- /dev/null +++ b/source/lib/common/md5sum.hpp @@ -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 +#include +#include +#include +#include +#include +#include + +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