AI/rocm-systems
2
0
Atdalīts 0
Kods Problēmas Izmaiņu pieprasījumi Darbības Pakotnes Projekti Laidieni Vikivietne Aktivitāte

GPU HW Counters via rocprofiler (#84)

Pārlūkot izejas kodu 
* Initial support for GPU hardware counters

* Update find modules for roctracer and rocprofiler

- /opt/rocm/{rocprofiler,roctracer} path is deprecated so tweak search procedure

* Improve ConfigCPack for MPI

* Update rocprofiler

- rocm_metrics()
- minor cleanup

* Update rocm find modules

* declare rocm_metrics + call in omnitrace-avail

* relocate omnitrace-launch-compiler

* REALPATH and find_modules

* Examples cmake (may drop)

* omnitrace-avail

- hw_counter categories
- init rocm

* setenv updates for rocprofiler in library.cpp and dl.cpp

* get_rocm_events config

* gpu::hip_device_count()

* rocm_metrics returns hardware_counters::info

* - relocated library/components/roctracer_callbacks.* to library/roctracer.*
- relocated library/components/rocprofiler.* to library/rocprofiler.*
- cleaned up rocprofiler.hpp
- added perfetto output of rocprofiler
- added timemory output of rocprofiler
- renamed omni.roctracer thread to roctracer.hip
- added roctracer.hsa thread name
- updated timemory submodule to support std::variant
- updated timemory submodule to support = in config value
- updated timemory submodule to support standalone storage
- updated timemory submodule to support new hw counter apis
- updated timemory submodule to prevent label/description caching in data_tracker

* update omnitrace-avail info_type generation

* Update timemory submodule

* rocprofiler component

* cmake formatting

* omnitrace-avail handle no GPUs

- Add -c command-line option for --categories
- support verbosity

* hsa_rsrc_factory throws exceptions

- throw exceptions to avoid aborting on HSA_STATUS_ERROR_NOT_INITIALIZED when advantageous
- removed duplicate specialization of is_available for component::rocprofiler

* rocprofiler symbols for when disabled

* Fix warning in omnitrace-avail

- std::stringstream from initializer list would use explicit constructor

* Fix finalization after settings are deleted

* Reorganized rocprofiler source

* Updated formatting

* Miscellaneous tweaks

- added using statements from timemory
- tweaked the main and thread bundle names
- fixed timemory header includes

[ROCm/rocprofiler-systems commit: 4208b5654c]
Šī revīzija ir iekļauta:
Jonathan R. Madsen
2022-07-17 21:52:09 -05:00
revīziju iesūtīja GitHub
vecāks aa7f2f63af
revīzija 7d1989f5a4
51 mainīti faili ar 3973 papildinājumiem un 390 dzēšanām
Lejupielādēt ielāpa failu Lejupielādēt izmaiņu failu Izvērst visus failus Savērst visus failus
projects/rocprofiler-systems/.github/workflows/formatting.yml
ārējs
+19 -1
Parādīt failu
@@ -74,7 +74,7 @@ jobs:
- name: clang-format
run: |
set +e
FILES=$(find include src examples -type f | egrep '\.hpp$|\.cpp$|\.cpp\.in$')
FILES=$(find source examples -type f | egrep '\.hpp$|\.cpp$|\.cpp\.in$')
FORMAT_OUT=$(clang-format-11 -output-replacements-xml ${FILES})
RET=$(echo ${FORMAT_OUT} | grep -c '<replacement ')
if [ "${RET}" -ne 0 ]; then
@@ -83,3 +83,21 @@ jobs:
git diff
exit ${RET}
fi
includes:
runs-on: ubuntu-20.04
steps:
- uses: actions/checkout@v2
- name: check-includes
run: |
set +e
FILES=$(find source examples -type f | egrep '\.hpp$|\.cpp$|\.cpp\.in$')
MATCHES=$(egrep 'include "timemory|include <bits' ${FILES})
if [ -n "${MATCHES}" ]; then
echo -e "\nError! Included timemory header with quotes or bits folder included\n"
echo -e "### MATCHES: ###"
echo -e "${MATCHES}"
echo -e "################"
exit 1
fi
projects/rocprofiler-systems/CMakeLists.txt
+26 -2
Parādīt failu
@@ -27,10 +27,28 @@ project(
DESCRIPTION "CPU/GPU Application tracing with static/dynamic binary instrumentation"
HOMEPAGE_URL "https://github.com/AMDResearch/omnitrace")
find_package(Git)
if(Git_FOUND AND EXISTS "${PROJECT_SOURCE_DIR}/.git")
execute_process(
COMMAND ${GIT_EXECUTABLE} describe --tags
OUTPUT_VARIABLE OMNITRACE_GIT_DESCRIBE
OUTPUT_STRIP_TRAILING_WHITESPACE)
execute_process(
COMMAND ${GIT_EXECUTABLE} rev-parse HEAD
OUTPUT_VARIABLE OMNITRACE_GIT_REVISION
OUTPUT_STRIP_TRAILING_WHITESPACE)
else()
set(OMNITRACE_GIT_DESCRIBE "v${OMNITRACE_VERSION}")
set(OMNITRACE_GIT_REVISION "")
endif()
message(
STATUS
"[${PROJECT_NAME}] version ${PROJECT_VERSION_MAJOR}.${PROJECT_VERSION_MINOR}.${PROJECT_VERSION_PATCH}"
"[${PROJECT_NAME}] version ${PROJECT_VERSION_MAJOR}.${PROJECT_VERSION_MINOR}.${PROJECT_VERSION_PATCH} (${FULL_VERSION_STRING})"
)
message(STATUS "[${PROJECT_NAME}] git revision: ${OMNITRACE_GIT_REVISION}")
message(STATUS "[${PROJECT_NAME}] git describe: ${OMNITRACE_GIT_DESCRIBE}")
set(CMAKE_MODULE_PATH ${PROJECT_SOURCE_DIR}/cmake ${PROJECT_SOURCE_DIR}/cmake/Modules
${PROJECT_SOURCE_DIR}/source/python/cmake ${CMAKE_MODULE_PATH})
set(BUILD_SHARED_LIBS
@@ -103,6 +121,8 @@ omnitrace_add_option(OMNITRACE_USE_HIP "Enable HIP support" ON)
omnitrace_add_option(OMNITRACE_USE_PAPI "Enable HW counter support via PAPI" ON)
omnitrace_add_option(OMNITRACE_USE_ROCTRACER "Enable roctracer support"
${OMNITRACE_USE_HIP})
omnitrace_add_option(OMNITRACE_USE_ROCPROFILER "Enable rocprofiler support"
${OMNITRACE_USE_HIP})
omnitrace_add_option(
OMNITRACE_USE_ROCM_SMI "Enable rocm-smi support for power/temp/etc. sampling"
${OMNITRACE_USE_HIP})
@@ -149,16 +169,20 @@ if(NOT OMNITRACE_USE_HIP)
set(OMNITRACE_USE_ROCTRACER
OFF
CACHE BOOL "Disabled via OMNITRACE_USE_HIP=OFF" FORCE)
set(OMNITRACE_USE_ROCPROFILER
OFF
CACHE BOOL "Disabled via OMNITRACE_USE_HIP=OFF" FORCE)
set(OMNITRACE_USE_ROCM_SMI
OFF
CACHE BOOL "Disabled via OMNITRACE_USE_HIP=OFF" FORCE)
elseif(
OMNITRACE_USE_HIP
AND NOT OMNITRACE_USE_ROCTRACER
AND NOT OMNITRACE_USE_ROCPROFILER
AND NOT OMNITRACE_USE_ROCM_SMI)
omnitrace_message(
AUTHOR_WARNING
"Setting OMNITRACE_USE_HIP=OFF because roctracer and rocm-smi options are disabled"
"Setting OMNITRACE_USE_HIP=OFF because roctracer, rocprofiler, and rocm-smi options are disabled"
)
set(OMNITRACE_USE_HIP OFF)
endif()
projects/rocprofiler-systems/VERSION
+1 -1
Parādīt failu
@@ -1 +1 @@
1.2.0
1.3.0dev0
projects/rocprofiler-systems/cmake/ConfigCPack.cmake
+4 -2
Parādīt failu
@@ -66,9 +66,9 @@ endif()
if(OMNITRACE_USE_MPI)
set(VALID_MPI_IMPLS "mpich" "openmpi")
if("${MPI_C_COMPILER_INCLUDE_DIRS}" MATCHES "openmpi")
if("${MPI_C_COMPILER_INCLUDE_DIRS};${MPI_C_HEADER_DIR}" MATCHES "openmpi")
set(OMNITRACE_MPI_IMPL "openmpi")
elseif("${MPI_C_COMPILER_INCLUDE_DIRS}" MATCHES "mpich")
elseif("${MPI_C_COMPILER_INCLUDE_DIRS};${MPI_C_HEADER_DIR}" MATCHES "mpich")
set(OMNITRACE_MPI_IMPL "mpich")
else()
message(
@@ -90,6 +90,8 @@ if(OMNITRACE_USE_MPI)
set(OMNITRACE_MPI_IMPL_UPPER "OpenMPI")
elseif("${OMNITRACE_MPI_IMPL}" STREQUAL "mpich")
set(OMNITRACE_MPI_IMPL_UPPER "MPICH")
else()
set(OMNITRACE_MPI_IMPL_UPPER "MPI")
endif()
set(OMNITRACE_CPACK_PACKAGE_SUFFIX
"${OMNITRACE_CPACK_PACKAGE_SUFFIX}-${OMNITRACE_MPI_IMPL_UPPER}")
projects/rocprofiler-systems/cmake/MacroUtilities.cmake
+5 -5
Parādīt failu
@@ -509,10 +509,10 @@ function(omnitrace_custom_compilation)
cmake_parse_arguments(COMP "GLOBAL;PROJECT" "COMPILER" "DIRECTORY;TARGET;SOURCE"
${ARGN})
# find omnitrace_launch_compiler
# find omnitrace-launch-compiler
find_program(
OMNITRACE_COMPILE_LAUNCHER
NAMES omnitrace_launch_compiler
NAMES omnitrace-launch-compiler
HINTS ${PROJECT_SOURCE_DIR} ${CMAKE_SOURCE_DIR}
PATHS ${PROJECT_SOURCE_DIR} ${CMAKE_SOURCE_DIR}
PATH_SUFFIXES scripts bin)
@@ -524,7 +524,7 @@ function(omnitrace_custom_compilation)
if(NOT OMNITRACE_COMPILE_LAUNCHER)
message(
FATAL_ERROR
"omnitrace could not find 'omnitrace_launch_compiler'. Please set '-DOMNITRACE_COMPILE_LAUNCHER=/path/to/launcher'"
"omnitrace could not find 'omnitrace-launch-compiler'. Please set '-DOMNITRACE_COMPILE_LAUNCHER=/path/to/launcher'"
)
endif()
@@ -797,7 +797,7 @@ function(OMNITRACE_INSTALL_TPL _TPL_TARGET _NEW_NAME _BUILD_TREE_DIR)
# build tree symbolic links
add_custom_target(
${_NEW_NAME}-library ALL
${CMAKE_COMMAND} -E create_symlink $<TARGET_FILE:${_TPL_TARGET}>
${CMAKE_COMMAND} -E copy_if_different $<TARGET_FILE:${_TPL_TARGET}>
${_TPL_PREFIX}${_NEW_NAME}${_TPL_SUFFIX}.${_TPL_VERSION}
COMMAND
${CMAKE_COMMAND} -E create_symlink
@@ -809,7 +809,7 @@ function(OMNITRACE_INSTALL_TPL _TPL_TARGET _NEW_NAME _BUILD_TREE_DIR)
${_BUILD_TREE_DIR}/${_TPL_PREFIX}${_NEW_NAME}${_TPL_SUFFIX}
WORKING_DIRECTORY ${_BUILD_TREE_DIR}
DEPENDS ${_TPL_TARGET}
COMMENT "Creating ${_NEW_NAME} symbolic links to ${_TPL_TARGET}...")
COMMENT "Creating ${_NEW_NAME} from ${_TPL_TARGET}...")
install(
FILES $<TARGET_FILE:${_TPL_TARGET}>
projects/rocprofiler-systems/cmake/Modules/FindROCmVersion.cmake
+9 -3
Parādīt failu
@@ -196,7 +196,7 @@ function(ROCM_VERSION_PARSE_VERSION_FILES)
if(EXISTS "${ROCmVersion_VERSION_FILE}" AND IS_ABSOLUTE
"${ROCmVersion_VERSION_FILE}")
get_filename_component(_VERSION_DIR "${ROCmVersion_VERSION_FILE}" PATH)
get_filename_component(_VERSION_DIR "${_VERSION_DIR}/.." ABSOLUTE)
get_filename_component(_VERSION_DIR "${_VERSION_DIR}/.." REALPATH)
set(ROCmVersion_DIR
"${_VERSION_DIR}"
CACHE PATH "Root path to ROCm's .info/${ROCmVersion_VERSION_FILE}"
@@ -227,8 +227,14 @@ function(ROCM_VERSION_PARSE_VERSION_FILES)
if(ROCmVersion_DIR)
set(_PATHS ${ROCmVersion_DIR})
else()
set(_PATHS ${ROCmVersion_DIR} ${ROCmVersion_ROOT} ${ROCmVersion_ROOT_DIR}
$ENV{CMAKE_PREFIX_PATH} ${CMAKE_PREFIX_PATH} ${ROCM_PATH} /opt/rocm)
set(_PATHS)
foreach(_DIR ${ROCmVersion_DIR} ${ROCmVersion_ROOT} ${ROCmVersion_ROOT_DIR}
$ENV{CMAKE_PREFIX_PATH} ${CMAKE_PREFIX_PATH} ${ROCM_PATH} /opt/rocm)
if(EXISTS ${_DIR})
get_filename_component(_ABS_DIR "${_DIR}" REALPATH)
list(APPEND _PATHS ${_ABS_DIR})
endif()
endforeach()
rocm_version_message(STATUS "ROCmVersion search paths: ${_PATHS}")
endif()
projects/rocprofiler-systems/cmake/Modules/Findrocm-smi.cmake
+18 -9
Parādīt failu
@@ -5,15 +5,24 @@ include(FindPackageHandleStandardArgs)
# ----------------------------------------------------------------------------------------#
# set(_ROCM_PATHS $ENV{ROCM_HOME} /opt/rocm)
if(NOT ROCM_PATH AND NOT "$ENV{ROCM_PATH}" STREQUAL "")
set(ROCM_PATH "$ENV{ROCM_PATH}")
endif()
foreach(_DIR ${ROCM_PATH} /opt/rocm /opt/rocm/rocm_smi)
if(EXISTS ${_DIR})
get_filename_component(_ABS_DIR "${_DIR}" REALPATH)
list(APPEND _ROCM_SMI_PATHS ${_ABS_DIR})
endif()
endforeach()
# ----------------------------------------------------------------------------------------#
find_path(
rocm-smi_ROOT_DIR
NAMES include/rocm_smi/rocm_smi.h
HINTS ${_ROCM_PATHS}
PATHS ${_ROCM_PATHS}
HINTS ${_ROCM_SMI_PATHS}
PATHS ${_ROCM_SMI_PATHS}
PATH_SUFFIXES rocm_smi)
mark_as_advanced(rocm-smi_ROOT_DIR)
@@ -23,9 +32,9 @@ mark_as_advanced(rocm-smi_ROOT_DIR)
find_path(
rocm-smi_INCLUDE_DIR
NAMES rocm_smi/rocm_smi.h
HINTS ${rocm-smi_ROOT_DIR} ${_ROCM_PATHS}
PATHS ${rocm-smi_ROOT_DIR} ${_ROCM_PATHS}
PATH_SUFFIXES rocm_smi/include rocm_smi)
HINTS ${rocm-smi_ROOT_DIR} ${_ROCM_SMI_PATHS}
PATHS ${rocm-smi_ROOT_DIR} ${_ROCM_SMI_PATHS}
PATH_SUFFIXES include rocm_smi/include)
mark_as_advanced(rocm-smi_INCLUDE_DIR)
@@ -34,8 +43,8 @@ mark_as_advanced(rocm-smi_INCLUDE_DIR)
find_library(
rocm-smi_LIBRARY
NAMES rocm_smi64 rocm_smi
HINTS ${rocm-smi_ROOT_DIR} ${_ROCM_PATHS}
PATHS ${rocm-smi_ROOT_DIR} ${_ROCM_PATHS}
HINTS ${rocm-smi_ROOT_DIR} ${_ROCM_SMI_PATHS}
PATHS ${rocm-smi_ROOT_DIR} ${_ROCM_SMI_PATHS}
PATH_SUFFIXES rocm_smi/lib rocm_smi/lib64 lib lib64)
if(rocm-smi_LIBRARY)
@@ -65,6 +74,6 @@ endif()
# ------------------------------------------------------------------------------#
unset(_ROCM_PATHS)
unset(_ROCM_SMI_PATHS)
# ------------------------------------------------------------------------------#
projects/rocprofiler-systems/cmake/Modules/Findrocprofiler.cmake
+100
Parādīt failu
@@ -0,0 +1,100 @@
# Distributed under the OSI-approved BSD 3-Clause License. See accompanying file
# Copyright.txt or https://cmake.org/licensing for details.
include(FindPackageHandleStandardArgs)
# ----------------------------------------------------------------------------------------#
if(NOT ROCM_PATH AND NOT "$ENV{ROCM_PATH}" STREQUAL "")
set(ROCM_PATH "$ENV{ROCM_PATH}")
endif()
foreach(_DIR ${ROCM_PATH} /opt/rocm /opt/rocm/rocprofiler)
if(EXISTS ${_DIR})
get_filename_component(_ABS_DIR "${_DIR}" REALPATH)
list(APPEND _ROCM_ROCPROFILER_PATHS ${_ABS_DIR})
endif()
endforeach()
# ----------------------------------------------------------------------------------------#
find_path(
rocprofiler_ROOT_DIR
NAMES include/rocprofiler/rocprofiler.h include/rocprofiler.h
HINTS ${_ROCM_ROCPROFILER_PATHS}
PATHS ${_ROCM_ROCPROFILER_PATHS}
PATH_SUFFIXES rocprofiler)
mark_as_advanced(rocprofiler_ROOT_DIR)
# ----------------------------------------------------------------------------------------#
find_path(
rocprofiler_INCLUDE_DIR
NAMES rocprofiler.h
HINTS ${rocprofiler_ROOT_DIR} ${_ROCM_ROCPROFILER_PATHS}
PATHS ${rocprofiler_ROOT_DIR} ${_ROCM_ROCPROFILER_PATHS}
PATH_SUFFIXES include include/rocprofiler rocprofiler/include)
mark_as_advanced(rocprofiler_INCLUDE_DIR)
find_path(
rocprofiler_hsa_INCLUDE_DIR
NAMES hsa.h
HINTS ${rocprofiler_ROOT_DIR} ${_ROCM_ROCPROFILER_PATHS}
PATHS ${rocprofiler_ROOT_DIR} ${_ROCM_ROCPROFILER_PATHS}
PATH_SUFFIXES include include/hsa)
mark_as_advanced(rocprofiler_hsa_INCLUDE_DIR)
# ----------------------------------------------------------------------------------------#
find_library(
rocprofiler_LIBRARY
NAMES rocprofiler64 rocprofiler
HINTS ${rocprofiler_ROOT_DIR}/rocprofiler ${rocprofiler_ROOT_DIR}
${_ROCM_ROCPROFILER_PATHS}
PATHS ${rocprofiler_ROOT_DIR}/rocprofiler ${rocprofiler_ROOT_DIR}
${_ROCM_ROCPROFILER_PATHS}
PATH_SUFFIXES lib lib64
NO_DEFAULT_PATH)
find_library(
rocprofiler_hsa-runtime_LIBRARY
NAMES hsa-runtime64 hsa-runtime
HINTS ${rocprofiler_ROOT_DIR} ${_ROCM_ROCPROFILER_PATHS}
PATHS ${rocprofiler_ROOT_DIR} ${_ROCM_ROCPROFILER_PATHS}
PATH_SUFFIXES lib lib64)
if(rocprofiler_LIBRARY)
get_filename_component(rocprofiler_LIBRARY_DIR "${rocprofiler_LIBRARY}" PATH CACHE)
endif()
mark_as_advanced(rocprofiler_LIBRARY rocprofiler_hsa-runtime_LIBRARY)
# ----------------------------------------------------------------------------------------#
find_package_handle_standard_args(
rocprofiler DEFAULT_MSG rocprofiler_ROOT_DIR rocprofiler_INCLUDE_DIR
rocprofiler_hsa_INCLUDE_DIR rocprofiler_LIBRARY rocprofiler_hsa-runtime_LIBRARY)
# ------------------------------------------------------------------------------#
if(rocprofiler_FOUND)
add_library(rocprofiler::rocprofiler INTERFACE IMPORTED)
add_library(rocprofiler::roctx INTERFACE IMPORTED)
set(rocprofiler_INCLUDE_DIRS ${rocprofiler_INCLUDE_DIR}
${rocprofiler_hsa_INCLUDE_DIR})
set(rocprofiler_LIBRARIES ${rocprofiler_LIBRARY} ${rocprofiler_hsa-runtime_LIBRARY})
set(rocprofiler_LIBRARY_DIRS ${rocprofiler_LIBRARY_DIR})
target_include_directories(
rocprofiler::rocprofiler INTERFACE ${rocprofiler_INCLUDE_DIR}
${rocprofiler_hsa_INCLUDE_DIR})
target_link_libraries(rocprofiler::rocprofiler INTERFACE ${rocprofiler_LIBRARIES})
endif()
# ------------------------------------------------------------------------------#
unset(_ROCM_ROCPROFILER_PATHS)
# ------------------------------------------------------------------------------#
projects/rocprofiler-systems/cmake/Modules/Findroctracer.cmake
+27 -18
Parādīt failu
@@ -5,15 +5,24 @@ include(FindPackageHandleStandardArgs)
# ----------------------------------------------------------------------------------------#
set(_ROCM_PATHS $ENV{ROCM_HOME} /opt/rocm /opt/rocm/roctracer)
if(NOT ROCM_PATH AND NOT "$ENV{ROCM_PATH}" STREQUAL "")
set(ROCM_PATH "$ENV{ROCM_PATH}")
endif()
foreach(_DIR ${ROCM_PATH} /opt/rocm /opt/rocm/roctracer)
if(EXISTS ${_DIR})
get_filename_component(_ABS_DIR "${_DIR}" REALPATH)
list(APPEND _ROCM_ROCTRACER_PATHS ${_ABS_DIR})
endif()
endforeach()
# ----------------------------------------------------------------------------------------#
find_path(
roctracer_ROOT_DIR
NAMES include/roctracer.h
HINTS ${_ROCM_PATHS}
PATHS ${_ROCM_PATHS}
NAMES include/roctracer/roctracer.h include/roctracer.h
HINTS ${_ROCM_ROCTRACER_PATHS}
PATHS ${_ROCM_ROCTRACER_PATHS}
PATH_SUFFIXES roctracer)
mark_as_advanced(roctracer_ROOT_DIR)
@@ -23,17 +32,17 @@ mark_as_advanced(roctracer_ROOT_DIR)
find_path(
roctracer_INCLUDE_DIR
NAMES roctracer.h
HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
PATH_SUFFIXES roctracer/include include)
HINTS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
PATH_SUFFIXES include include/roctracer roctracer/include)
mark_as_advanced(roctracer_INCLUDE_DIR)
find_path(
roctracer_hsa_INCLUDE_DIR
NAMES hsa.h
HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
HINTS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
PATH_SUFFIXES include include/hsa)
mark_as_advanced(roctracer_hsa_INCLUDE_DIR)
@@ -43,22 +52,22 @@ mark_as_advanced(roctracer_hsa_INCLUDE_DIR)
find_library(
roctracer_LIBRARY
NAMES roctracer64 roctracer
HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
HINTS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
PATH_SUFFIXES lib lib64)
find_library(
roctracer_roctx_LIBRARY
NAMES roctx64 roctx
HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
HINTS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
PATH_SUFFIXES lib lib64)
find_library(
roctracer_kfdwrapper_LIBRARY
NAMES kfdwrapper64 kfdwrapper
HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
HINTS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
PATH_SUFFIXES lib lib64)
find_package(hsakmt)
@@ -70,8 +79,8 @@ else()
find_library(
roctracer_hsakmt_LIBRARY
NAMES hsakmt
HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
HINTS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_ROCTRACER_PATHS}
PATH_SUFFIXES lib lib64)
endif()
@@ -122,6 +131,6 @@ endif()
# ------------------------------------------------------------------------------#
unset(_ROCM_PATHS)
unset(_ROCM_ROCTRACER_PATHS)
# ------------------------------------------------------------------------------#
projects/rocprofiler-systems/cmake/Packages.cmake
+37 -7
Parādīt failu
@@ -16,6 +16,8 @@ omnitrace_add_interface_library(
omnitrace_add_interface_library(omnitrace-hip "Provides flags and libraries for HIP")
omnitrace_add_interface_library(omnitrace-roctracer
"Provides flags and libraries for roctracer")
omnitrace_add_interface_library(omnitrace-rocprofiler
"Provides flags and libraries for rocprofiler")
omnitrace_add_interface_library(omnitrace-rocm-smi
"Provides flags and libraries for rocm-smi")
omnitrace_add_interface_library(omnitrace-mpi "Provides MPI or MPI headers")
@@ -31,9 +33,15 @@ omnitrace_add_interface_library(omnitrace-compile-definitions "Compile definitio
# libraries with relevant compile definitions
set(OMNITRACE_EXTENSION_LIBRARIES
omnitrace::omnitrace-hip omnitrace::omnitrace-roctracer omnitrace::omnitrace-rocm-smi
omnitrace::omnitrace-mpi omnitrace::omnitrace-ptl omnitrace::omnitrace-ompt
omnitrace::omnitrace-papi omnitrace::omnitrace-perfetto)
omnitrace::omnitrace-hip
omnitrace::omnitrace-roctracer
omnitrace::omnitrace-rocprofiler
omnitrace::omnitrace-rocm-smi
omnitrace::omnitrace-mpi
omnitrace::omnitrace-ptl
omnitrace::omnitrace-ompt
omnitrace::omnitrace-papi
omnitrace::omnitrace-perfetto)
target_include_directories(
omnitrace-headers INTERFACE ${PROJECT_SOURCE_DIR}/source/lib/omnitrace
@@ -45,6 +53,14 @@ target_link_libraries(omnitrace-headers INTERFACE omnitrace::omnitrace-threading
# ensure the env overrides the appending /opt/rocm later
string(REPLACE ":" ";" CMAKE_PREFIX_PATH "$ENV{CMAKE_PREFIX_PATH};${CMAKE_PREFIX_PATH}")
set(OMNITRACE_DEFAULT_ROCM_PATH
/opt/rocm
CACHE PATH "Default search path for ROCM")
if(EXISTS ${OMNITRACE_DEFAULT_ROCM_PATH})
get_filename_component(OMNITRACE_DEFAULT_ROCM_PATH "${OMNITRACE_DEFAULT_ROCM_PATH}"
REALPATH)
endif()
# ----------------------------------------------------------------------------------------#
#
# Threading
@@ -84,7 +100,7 @@ endforeach()
# ----------------------------------------------------------------------------------------#
if(OMNITRACE_USE_HIP)
list(APPEND CMAKE_PREFIX_PATH /opt/rocm)
list(APPEND CMAKE_PREFIX_PATH ${OMNITRACE_DEFAULT_ROCM_PATH})
find_package(hip ${omnitrace_FIND_QUIETLY} REQUIRED)
omnitrace_target_compile_definitions(omnitrace-hip INTERFACE OMNITRACE_USE_HIP)
target_link_libraries(omnitrace-hip INTERFACE hip::host)
@@ -97,7 +113,7 @@ endif()
# ----------------------------------------------------------------------------------------#
if(OMNITRACE_USE_ROCTRACER)
list(APPEND CMAKE_PREFIX_PATH /opt/rocm)
list(APPEND CMAKE_PREFIX_PATH ${OMNITRACE_DEFAULT_ROCM_PATH})
find_package(roctracer ${omnitrace_FIND_QUIETLY} REQUIRED)
omnitrace_target_compile_definitions(omnitrace-roctracer
INTERFACE OMNITRACE_USE_ROCTRACER)
@@ -106,6 +122,20 @@ if(OMNITRACE_USE_ROCTRACER)
set(CMAKE_INSTALL_RPATH "${CMAKE_INSTALL_RPATH}:${roctracer_LIBRARY_DIRS}")
endif()
# ----------------------------------------------------------------------------------------#
#
# rocprofiler
#
# ----------------------------------------------------------------------------------------#
if(OMNITRACE_USE_ROCPROFILER)
list(APPEND CMAKE_PREFIX_PATH ${OMNITRACE_DEFAULT_ROCM_PATH})
find_package(rocprofiler ${omnitrace_FIND_QUIETLY} REQUIRED)
omnitrace_target_compile_definitions(omnitrace-rocprofiler
INTERFACE OMNITRACE_USE_ROCPROFILER)
target_link_libraries(omnitrace-rocprofiler INTERFACE rocprofiler::rocprofiler)
set(CMAKE_INSTALL_RPATH "${CMAKE_INSTALL_RPATH}:${rocprofiler_LIBRARY_DIRS}")
endif()
# ----------------------------------------------------------------------------------------#
#
# rocm-smi
@@ -113,7 +143,7 @@ endif()
# ----------------------------------------------------------------------------------------#
if(OMNITRACE_USE_ROCM_SMI)
list(APPEND CMAKE_PREFIX_PATH /opt/rocm)
list(APPEND CMAKE_PREFIX_PATH ${OMNITRACE_DEFAULT_ROCM_PATH})
find_package(rocm-smi ${omnitrace_FIND_QUIETLY} REQUIRED)
omnitrace_target_compile_definitions(omnitrace-rocm-smi
INTERFACE OMNITRACE_USE_ROCM_SMI)
@@ -545,7 +575,7 @@ omnitrace_checkout_git_submodule(
RELATIVE_PATH external/timemory
WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}
REPO_URL https://github.com/NERSC/timemory.git
REPO_BRANCH gpu-kernel-instrumentation)
REPO_BRANCH omnitrace)
omnitrace_save_variables(
BUILD_CONFIG VARIABLES BUILD_SHARED_LIBS BUILD_STATIC_LIBS
projects/rocprofiler-systems/examples/lulesh/CMakeLists.txt
+2
Parādīt failu
@@ -2,6 +2,8 @@ cmake_minimum_required(VERSION 3.16 FATAL_ERROR)
project(lulesh LANGUAGES C CXX)
set(CMAKE_BUILD_TYPE "RelWithDebInfo")
list(INSERT CMAKE_MODULE_PATH 0 ${PROJECT_SOURCE_DIR}/cmake/Modules)
option(LULESH_BUILD_KOKKOS "Build Kokkos from submodule" ON)
projects/rocprofiler-systems/examples/openmp/CMakeLists.txt
+4 -4
Parādīt failu
@@ -23,10 +23,7 @@ else()
AND LIBOMP_LIBRARY
AND COMMAND omnitrace_custom_compilation)
target_compile_options(openmp-common PUBLIC -W -Wall -fopenmp=libomp)
target_compile_options(openmp-cg PRIVATE -W -Wall -fopenmp=libomp)
target_link_libraries(openmp-cg PRIVATE ${LIBOMP_LIBRARY})
target_compile_options(openmp-lu PRIVATE -W -Wall -fopenmp=libomp)
target_link_libraries(openmp-lu PRIVATE ${LIBOMP_LIBRARY})
target_link_libraries(openmp-common PUBLIC ${LIBOMP_LIBRARY})
omnitrace_custom_compilation(COMPILER ${CLANGXX_EXECUTABLE} TARGET openmp-common)
omnitrace_custom_compilation(COMPILER ${CLANGXX_EXECUTABLE} TARGET openmp-cg)
omnitrace_custom_compilation(COMPILER ${CLANGXX_EXECUTABLE} TARGET openmp-lu)
@@ -36,6 +33,9 @@ else()
endif()
endif()
target_link_libraries(openmp-cg PRIVATE openmp-common)
target_link_libraries(openmp-lu PRIVATE openmp-common)
if(OMNITRACE_INSTALL_EXAMPLES)
install(
TARGETS openmp-cg openmp-lu
projects/rocprofiler-systems/examples/transpose/transpose.cpp
+19 -17
Parādīt failu
@@ -93,13 +93,18 @@ transpose_a(int* in, int* out, int M, int N)
out[idx] = tile[threadIdx.x][threadIdx.y];
}
namespace
{
size_t nthreads = 2;
size_t nitr = 500;
size_t nsync = 10;
} // namespace
void
run(int rank, int tid, hipStream_t stream, int argc, char** argv)
{
size_t nitr = 500;
size_t nsync = 10;
unsigned int M = 4960 * 2;
unsigned int N = 4960 * 2;
unsigned int M = 4960 * 2;
unsigned int N = 4960 * 2;
if(argc > 2) nitr = atoll(argv[2]);
if(argc > 3) nsync = atoll(argv[3]);
@@ -186,29 +191,26 @@ do_a2a(int rank)
int
main(int argc, char** argv)
{
int rank = 0;
int size = 1;
int nthreads = 2;
int nitr = 5000;
size_t nsync = 10;
int rank = 0;
int size = 1;
for(int i = 1; i < argc; ++i)
{
auto _arg = std::string{ argv[i] };
if(_arg == "?" || _arg == "-h" || _arg == "--help")
{
fprintf(stderr,
"usage: transpose [NUM_THREADS (%i)] [NUM_ITERATION (%i)] "
"usage: transpose [NUM_THREADS (%zu)] [NUM_ITERATION (%zu)] "
"[SYNC_EVERY_N_ITERATIONS (%zu)]\n",
nthreads, nitr, nsync);
exit(EXIT_SUCCESS);
}
}
if(argc > 1) nthreads = atoi(argv[1]);
if(argc > 2) nitr = atoi(argv[2]);
if(argc > 1) nthreads = atoll(argv[1]);
if(argc > 2) nitr = atoll(argv[2]);
if(argc > 3) nsync = atoll(argv[3]);
printf("[transpose] Number of threads: %i\n", nthreads);
printf("[transpose] Number of iterations: %i\n", nitr);
printf("[transpose] Number of threads: %zu\n", nthreads);
printf("[transpose] Number of iterations: %zu\n", nitr);
printf("[transpose] Syncing every %zu iterations\n", nsync);
#if defined(USE_MPI)
@@ -233,14 +235,14 @@ main(int argc, char** argv)
{
std::vector<std::thread> _threads{};
std::vector<hipStream_t> _streams(nthreads);
for(int i = 0; i < nthreads; ++i)
for(size_t i = 0; i < nthreads; ++i)
HIP_API_CALL(hipStreamCreate(&_streams.at(i)));
for(int i = 1; i < nthreads; ++i)
for(size_t i = 1; i < nthreads; ++i)
_threads.emplace_back(run, rank, i, _streams.at(i), argc, argv);
run(rank, 0, _streams.at(0), argc, argv);
for(auto& itr : _threads)
itr.join();
for(int i = 0; i < nthreads; ++i)
for(size_t i = 0; i < nthreads; ++i)
HIP_API_CALL(hipStreamDestroy(_streams.at(i)));
}
HIP_API_CALL(hipDeviceSynchronize());
projects/rocprofiler-systems/external/timemory
ārējs
+1 -1
Submodule projects/rocprofiler-systems/external/timemory updated: 33ef109c08...1c03bec283
projects/rocprofiler-systems/scripts/gperftools-cpu-profile.sh
+42 -61
Parādīt failu
@@ -1,16 +1,18 @@
#!/bin/bash
EXE=$(basename ${1})
DIR=cpu.prof.${EXE}
DIR=gperftools-output
mkdir -p ${DIR}
# gperf settings
# gperftools settings
: ${N:=0}
: ${GPERF_PROFILE:=""}
: ${GPERF_PROFILE_BASE:=${DIR}/gperf}
: ${GPERFTOOLS_PROFILE:=""}
: ${GPERFTOOLS_PROFILE_BASE:=${DIR}/prof.${EXE}}
: ${MALLOCSTATS:=1}
: ${CPUPROFILE_FREQUENCY:=250}
: ${CPUPROFILE_REALTIME:=1}
: ${PPROF:=$(which google-pprof)}
: ${PPROF:=$(which pprof)}
# rendering settings
: ${INTERACTIVE:=0}
@@ -24,17 +26,18 @@ if [ "$(uname)" = "Darwin" ]; then
IMG_FORMAT="jpg"
fi
fi
run-verbose()
{
echo "${@}" 1>&2
echo "### ${@} ###" 1>&2
eval ${@}
}
while [ -z "${GPERF_PROFILE}" ]
while [ -z "${GPERFTOOLS_PROFILE}" ]
do
TEST_FILE=${GPERF_PROFILE_BASE}.${N}
TEST_FILE=${GPERFTOOLS_PROFILE_BASE}.${N}
if [ ! -f "${TEST_FILE}" ]; then
GPERF_PROFILE=${TEST_FILE}
GPERFTOOLS_PROFILE=${TEST_FILE}
fi
N=$((${N}+1))
done
@@ -43,10 +46,10 @@ export MALLOCSTATS
export CPUPROFILE_FREQUENCY
export CPUPROFILE_REALTIME
echo -e "\n\t--> Outputting profile to '${GPERF_PROFILE}'...\n"
echo -e "\n\t--> Outputting profile to '${GPERFTOOLS_PROFILE}'...\n"
# remove profile file if unsucessful execution
cleanup-failure() { set +v ; echo "failure"; rm -f ${GPERF_PROFILE}; exit 1; }
cleanup-failure() { set +v ; echo "failure"; rm -f ${GPERFTOOLS_PROFILE}; exit 1; }
trap cleanup-failure SIGHUP SIGINT SIGQUIT SIGILL SIGABRT SIGKILL
ADD_LIBS()
@@ -73,66 +76,44 @@ ADD_PRELOAD()
done
}
run-pprof()
{
if [ -n "${PPROF}" ]; then
run-verbose ${PPROF} ${ADD_LIB_LIST} ${PPROF_ARGS} "${@}"
else
echo -e "neither google-pprof nor pprof were found!"
exit 1
fi
}
# configure pre-loading of profiler library
PROJECT_LIBRARIES="$(find ${PWD} -type f | egrep 'libtimemory|libctimemory' | egrep -v '\.a$|\.dSYM' | egrep '\.so$|\.dylib$')"
run-verbose ADD_LIBS ${PROJECT_LIBRARIES}
if [ "$(uname)" = "Darwin" ]; then
run-verbose ADD_PRELOAD $(otool -L ${1} | egrep 'profiler' | awk '{print $1}')
LIBS=$(echo ${LIBS} | sed 's/^://g')
if [ -n "${LIBS}" ]; then
export DYLD_FORCE_FLAT_NAMESPACE=1
echo "DYLD_INSERT_LIBRARIES=${LIBS}"
fi
else
run-verbose ADD_PRELOAD $(ldd ${1} | egrep 'profiler' | awk '{print $(NF-1)}') /usr/lib/x86_64-linux-gnu/libprofiler.so
LIBS=$(echo ${LIBS} | sed 's/^://g')
if [ -n "${LIB}" ]; then
echo "LD_PRELOAD=${LIBS}"
fi
fi
for i in $(find ${PWD} -type f | egrep 'libomnitrace' | egrep -v '\.a$' | egrep '\.so$') $(ldd ${1} | awk '{print $(NF-1)}')
do
if [ -f "${i}" ]; then run-verbose ADD_LIBS "${i}"; fi
done
run-verbose ADD_PRELOAD $(ldd ${1} | egrep 'profiler' | awk '{print $(NF-1)}') /usr/lib/$(uname -m)-linux-gnu/libprofiler.so
LIBS=$(echo ${LIBS} | sed 's/^://g')
set -e
# run the application
if [ "$(uname)" = "Darwin" ]; then
eval DYLD_INSERT_LIBRARIES=${LIBS} CPUPROFILE_FREQUENCY=${CPUPROFILE_FREQUENCY} CPUPROFILE=${GPERF_PROFILE} $@ | tee ${GPERF_PROFILE}.log
else
eval LD_PRELOAD=${LIBS} CPUPROFILE_FREQUENCY=${CPUPROFILE_FREQUENCY} CPUPROFILE=${GPERF_PROFILE} $@ | tee ${GPERF_PROFILE}.log
fi
LD_PRELOAD=${LIBS} CPUPROFILE_FREQUENCY=${CPUPROFILE_FREQUENCY} CPUPROFILE=${GPERFTOOLS_PROFILE} ${@} | tee ${GPERFTOOLS_PROFILE}.log
set +e
echo-dart-measurement()
{
local _NAME=${1}
local _TYPE=${2}
local _PATH=${3}
echo "<DartMeasurementFile name=\"${_NAME}\" type=\"image/${_TYPE}\">${_PATH}</DartMeasurementFile>"
}
# generate the results
EXT=so
if [ "$(uname)" = "Darwin" ]; then EXT=dylib; fi
if [ -f "${GPERF_PROFILE}" ]; then
: ${PPROF:=$(which google-pprof)}
: ${PPROF:=$(which pprof)}
if [ -n "${PPROF}" ]; then
run-verbose ${PPROF} --text ${ADD_LIB_LIST} ${PPROF_ARGS} ${1} ${GPERF_PROFILE} 1> ${GPERF_PROFILE}.txt.tmp
run-verbose cat ${GPERF_PROFILE}.txt.tmp | c++filt -n -t &> ${GPERF_PROFILE}.txt
run-verbose ${PPROF} --text --cum ${ADD_LIB_LIST} ${PPROF_ARGS} ${1} ${GPERF_PROFILE} 1> ${GPERF_PROFILE}.cum.txt.tmp
run-verbose cat ${GPERF_PROFILE}.cum.txt.tmp | c++filt -n -t &> ${GPERF_PROFILE}.cum.txt
rm -f *.txt.tmp
# if dot is available
if [ -n "$(which dot)" ]; then
run-verbose ${PPROF} --dot ${ADD_LIB_LIST} ${PPROF_ARGS} ${1} ${GPERF_PROFILE} 1> ${GPERF_PROFILE}.dot
run-verbose dot ${DOT_ARGS} -T${IMG_FORMAT} ${GPERF_PROFILE}.dot -o ${GPERF_PROFILE}.${IMG_FORMAT}
echo-dart-measurement ${GPERF_PROFILE}.${IMG_FORMAT} ${IMG_FORMAT} ${PWD}/${GPERF_PROFILE}.${IMG_FORMAT}
fi
if [ "${INTERACTIVE}" -gt 0 ]; then
run-verbose ${PPROF} ${ADD_LIB_LIST} ${PPROF_ARGS} ${1} ${GPERF_PROFILE}
fi
else
echo -e "google-pprof/pprof not found!"
if [ -f "${GPERFTOOLS_PROFILE}" ]; then
run-pprof --text ${1} ${GPERFTOOLS_PROFILE} | c++filt -n -t 1> ${GPERFTOOLS_PROFILE}.txt
run-pprof ${PPROF} --text --cum ${1} ${GPERFTOOLS_PROFILE} | c++filt -n -t 1> ${GPERFTOOLS_PROFILE}.cum.txt
# if dot is available
if [ -n "$(which dot)" ]; then
run-pprof ${PPROF} --dot ${1} ${GPERFTOOLS_PROFILE} 1> ${GPERFTOOLS_PROFILE}.dot
run-verbose $(which dot) ${DOT_ARGS} -T${IMG_FORMAT} ${GPERFTOOLS_PROFILE}.dot -o ${GPERFTOOLS_PROFILE}.${IMG_FORMAT}
fi
if [ "${INTERACTIVE}" -gt 0 ]; then
run-pprof ${PPROF} ${1} ${GPERFTOOLS_PROFILE}
fi
else
echo -e "profile file \"${GPERF_PROFILE}\" not found!"
echo -e "profile file \"${GPERFTOOLS_PROFILE}\" not found!"
ls -la
exit 1
fi
projects/rocprofiler-systems/scripts/omnitrace_launch_compiler → projects/rocprofiler-systems/scripts/omnitrace-launch-compiler
Parādīt failu
projects/rocprofiler-systems/source/bin/omnitrace-avail/avail.cpp
+58 -10
Parādīt failu
@@ -30,6 +30,8 @@
#include "info_type.hpp"
#include "library/config.hpp"
#include "library/gpu.hpp"
#include "library/rocprofiler.hpp"
#include <timemory/components.hpp>
#include <timemory/components/definition.hpp>
@@ -57,6 +59,12 @@
#include <utility>
#include <vector>
#if defined(OMNITRACE_USE_HIP) && OMNITRACE_USE_HIP > 0
# include <hip/hip_runtime.h>
#elif !defined(OMNITRACE_USE_HIP)
# define OMNITRACE_USE_HIP 0
#endif
#if defined(TIMEMORY_UNIX)
# include <sys/ioctl.h> // ioctl() and TIOCGWINSZ
# include <unistd.h> // for STDOUT_FILENO
@@ -98,6 +106,8 @@ void
write_hw_counter_info(std::ostream&, const array_t<bool, N>& = {},
const array_t<bool, N>& = {}, const array_t<string_t, N>& = {});
int gpu_count = 0;
//--------------------------------------------------------------------------------------//
int
@@ -125,6 +135,8 @@ main(int argc, char** argv)
}
}
}
_category_options.emplace("hw_counters::CPU");
_category_options.emplace("hw_counters::GPU");
array_t<bool, TOTAL> options = { false, false, false, false, false, false, false };
array_t<string_t, TOTAL> fields = {};
@@ -164,6 +176,11 @@ main(int argc, char** argv)
parser.add_argument({ "--debug" }, "Enable debug messages")
.max_count(1)
.action([](parser_t& p) { debug_msg = p.get<bool>("debug"); });
parser.add_argument({ "--verbose" }, "Enable informational messages")
.max_count(1)
.action([](parser_t& p) {
verbose_level = (p.get_count("verbose") == 0) ? 1 : p.get<int>("verbose");
});
parser.add_argument({ "-a", "--all" }, "Print all available info")
.max_count(1)
.action([&](parser_t& p) {
@@ -297,7 +314,7 @@ main(int argc, char** argv)
"Display the output filename for the component")
.max_count(1);
parser
.add_argument({ "--categories" },
.add_argument({ "-c", "--categories" },
"Display the category information (use --list-categories to see "
"the available categories)")
.dtype("string")
@@ -399,6 +416,20 @@ main(int argc, char** argv)
return EXIT_FAILURE;
}
#if OMNITRACE_USE_HIP > 0
// initialize HIP and call rocm_metrics() which add choices to OMNITRACE_ROCM_EVENTS
// setting
auto _status = hipGetDeviceCount(&gpu_count);
if(gpu_count > 0 && _status == hipSuccess)
{
(void) omnitrace::rocprofiler::rocm_metrics();
}
else
{
verbprintf(0, "No HIP devices found. GPU HW counters will not be available\n");
}
#endif
auto _parser_set_if_exists = [&parser](auto& _var, const std::string& _opt) {
using Tp = decay_t<decltype(_var)>;
if(parser.exists(_opt)) _var = parser.get<Tp>(_opt);
@@ -923,10 +954,13 @@ write_hw_counter_info(std::ostream& os, const array_t<bool, N>& options,
static_assert(N >= num_hw_counter_options,
"Error! Too few hw counter options + fields");
using width_type = array_t<int64_t, N>;
using width_bool = array_t<bool, N>;
using width_type = array_t<int64_t, N>;
using width_bool = array_t<bool, N>;
using hwcounter_info_t = std::vector<tim::hardware_counters::info>;
auto _papi_events = tim::papi::available_events_info();
auto _rocm_events =
(gpu_count > 0) ? omnitrace::rocprofiler::rocm_metrics() : hwcounter_info_t{};
auto _process_counters = [](auto& _events, int32_t _offset) {
for(auto& itr : _events)
@@ -939,13 +973,25 @@ write_hw_counter_info(std::ostream& os, const array_t<bool, N>& options,
int32_t _offset = 0;
_offset += _process_counters(_papi_events, _offset);
_offset += _process_counters(_rocm_events, _offset);
using hwcounter_info_t = std::vector<tim::hardware_counters::info>;
auto fields = std::vector<hwcounter_info_t>{ _papi_events };
auto subcategories = std::vector<std::string>{ "CPU", "GPU", "" };
array_t<string_t, N> _labels = { "HARDWARE COUNTER", "AVAILABLE", "SUMMARY",
auto fields = std::vector<hwcounter_info_t>{ _papi_events, _rocm_events };
auto subcategories = std::vector<std::string>{ "CPU", "GPU", "" };
array_t<string_t, N> _labels = { "HARDWARE COUNTER", "AVAILABLE", "SUMMARY",
"DESCRIPTION" };
array_t<bool, N> _center = { false, true, false, false };
array_t<bool, N> _center = { false, true, false, false };
for(size_t i = 0; i < subcategories.size(); ++i)
{
if(i >= fields.size()) break;
if(!category_view.empty() && category_view.count(subcategories.at(i)) == 0 &&
category_view.count(std::string{ "hw_counters::" } + subcategories.at(i)) == 0)
fields.at(i).clear();
if(!is_category_selected(subcategories.at(i)) &&
!is_category_selected(std::string{ "hw_counters::" } + subcategories.at(i)))
fields.at(i).clear();
if(fields.at(i).empty()) subcategories.at(i).clear();
}
width_type _widths;
width_bool _wusing;
@@ -986,14 +1032,15 @@ write_hw_counter_info(std::ostream& os, const array_t<bool, N>& options,
os << "\n" << banner(_widths, _wusing, '-');
size_t nitr = 0;
size_t nout = 0;
for(const auto& fitr : fields)
{
auto idx = nitr++;
if(idx < subcategories.size())
{
if(!markdown && idx != 0) os << banner(_widths, _wusing, '-');
if(subcategories.at(idx).length() > 0)
if(!markdown && nout != 0) os << banner(_widths, _wusing, '-');
if(!subcategories.at(idx).empty())
{
os << global_delim;
if(options[0])
@@ -1008,6 +1055,7 @@ write_hw_counter_info(std::ostream& os, const array_t<bool, N>& options,
}
os << "\n";
if(!markdown) os << banner(_widths, _wusing, '-');
++nout;
}
}
else
projects/rocprofiler-systems/source/bin/omnitrace-avail/common.cpp
+7 -3
Parādīt failu
@@ -38,7 +38,6 @@ bool alphabetical = false;
bool available_only = false;
bool all_info = false;
bool force_brief = false;
bool debug_msg = false;
bool case_insensitive = false;
bool regex_hl = false;
bool expand_keys = false;
@@ -50,7 +49,11 @@ int32_t padding = 4;
str_vec_t regex_keys = {};
str_vec_t category_regex_keys = {};
str_set_t category_view = {};
std::stringstream lerr = {};
std::stringstream lerr{};
bool debug_msg = tim::get_env<bool>("OMNITRACE_DEBUG_AVAIL", settings::debug());
int32_t verbose_level =
tim::get_env<int32_t>("OMNITRACE_VERBOSE_AVAIL", settings::verbose());
// explicit setting names to exclude
std::set<std::string> settings_exclude = {
@@ -321,7 +324,8 @@ process_categories(parser_t& p, const str_set_t& _category_options)
if(_category_options.count(itr) == 0)
{
if(!_is_shorthand("component") && !_is_shorthand("settings"))
if(!_is_shorthand("component") && !_is_shorthand("settings") &&
!_is_shorthand("hw_counters"))
throw std::runtime_error(
itr + " is not a valid category. Use --list-categories to view "
"valid categories");
projects/rocprofiler-systems/source/bin/omnitrace-avail/common.hpp
+34
Parādīt failu
@@ -99,6 +99,7 @@ extern int32_t max_width;
extern int32_t num_cols;
extern int32_t min_width;
extern int32_t padding;
extern int32_t verbose_level;
extern str_vec_t regex_keys;
extern str_vec_t category_regex_keys;
extern str_set_t category_view;
@@ -143,3 +144,36 @@ remove(std::string inp, const std::set<std::string>& entries);
bool
file_exists(const std::string&);
// control debug printf statements
#define errprintf(LEVEL, ...) \
{ \
if(werror || LEVEL < 0) \
{ \
if(debug_msg || verbose_level >= LEVEL) \
fprintf(stderr, "[omnitrace][avail] Error! " __VA_ARGS__); \
char _buff[FUNCNAMELEN]; \
sprintf(_buff, "[omnitrace][avail] Error! " __VA_ARGS__); \
throw std::runtime_error(std::string{ _buff }); \
} \
else \
{ \
if(debug_msg || verbose_level >= LEVEL) \
fprintf(stderr, "[omnitrace][avail] Warning! " __VA_ARGS__); \
} \
fflush(stderr); \
}
// control verbose printf statements
#define verbprintf(LEVEL, ...) \
{ \
if(debug_msg || verbose_level >= LEVEL) \
fprintf(stderr, "[omnitrace][avail] " __VA_ARGS__); \
fflush(stderr); \
}
#define verbprintf_bare(LEVEL, ...) \
{ \
if(debug_msg || verbose_level >= LEVEL) fprintf(stderr, __VA_ARGS__); \
fflush(stderr); \
}
projects/rocprofiler-systems/source/bin/omnitrace-avail/generate_config.cpp
+4 -3
Parādīt failu
@@ -230,7 +230,7 @@ generate_config(std::string _config_file, const std::set<std::string>& _config_f
_ar->startNode();
(*_ar)(cereal::make_nvp("version", std::string{ OMNITRACE_VERSION_STRING }));
(*_ar)(cereal::make_nvp("date", tim::get_local_datetime("%F_%H.%M", &_time)));
settings::serialize_settings(*_ar);
settings::serialize_settings(*_ar, *_settings);
_ar->finishNode();
};
@@ -311,6 +311,7 @@ generate_config(std::string _config_file, const std::set<std::string>& _config_f
});
else
{
_settings->ordering();
std::sort(_data.begin(), _data.end(), [](auto _lhs, auto _rhs) {
auto _lomni = _lhs->get_categories().count("omnitrace") > 0;
auto _romni = _rhs->get_categories().count("omnitrace") > 0;
@@ -419,7 +420,7 @@ update_choices(std::shared_ptr<settings> _settings)
{
std::vector<info_type> _info = get_component_info<TIMEMORY_NATIVE_COMPONENTS_END>();
if(settings::verbose() >= 2 || settings::debug())
if(_settings->get_verbose() >= 2 || _settings->get_debug())
printf("[omnitrace-avail] # of component found: %zu\n", _info.size());
_info.erase(std::remove_if(_info.begin(), _info.end(),
@@ -442,7 +443,7 @@ update_choices(std::shared_ptr<settings> _settings)
_component_choices.reserve(_info.size());
for(const auto& itr : _info)
_component_choices.emplace_back(itr.id_type());
if(settings::verbose() >= 2 || settings::debug())
if(_settings->get_verbose() >= 2 || _settings->get_debug())
printf("[omnitrace-avail] # of component choices: %zu\n",
_component_choices.size());
_settings->find("OMNITRACE_TIMEMORY_COMPONENTS")
projects/rocprofiler-systems/source/bin/omnitrace-avail/info_type.cpp
+1
Parādīt failu
@@ -30,6 +30,7 @@
#include "library/components/mpi_gotcha.hpp"
#include "library/components/omnitrace.hpp"
#include "library/components/pthread_gotcha.hpp"
#include "library/components/rocprofiler.hpp"
#include "library/components/roctracer.hpp"
#include "library/components/user_region.hpp"
projects/rocprofiler-systems/source/lib/common/defines.h
+1
Parādīt failu
@@ -31,6 +31,7 @@
#define OMNITRACE_HOT OMNITRACE_ATTRIBUTE(hot)
#define OMNITRACE_CONST OMNITRACE_ATTRIBUTE(const)
#define OMNITRACE_PURE OMNITRACE_ATTRIBUTE(pure)
#define OMNITRACE_WEAK OMNITRACE_ATTRIBUTE(weak)
#define OMNITRACE_PACKED OMNITRACE_ATTRIBUTE(__packed__)
#define OMNITRACE_PACKED_ALIGN(VAL) OMNITRACE_PACKED OMNITRACE_ATTRIBUTE(__aligned__(VAL))
projects/rocprofiler-systems/source/lib/omnitrace-dl/dl.cpp
+19
Parādīt failu
@@ -157,6 +157,25 @@ struct OMNITRACE_HIDDEN_API indirect
setenv("HSA_TOOLS_LIB", _omni_hsa_lib.c_str(), 0);
#endif
#if defined(OMNITRACE_USE_ROCPROFILER) && OMNITRACE_USE_ROCPROFILER > 0
auto _rocm_path = get_env("ROCM_PATH", "/opt/rocm");
auto _rocp_metrics = common::join('/', _rocm_path, "rocprofiler/lib/metrics.xml");
setenv("HSA_TOOLS_LIB", m_omnilib.c_str(), 0);
setenv("ROCP_TOOL_LIB", m_omnilib.c_str(), 0);
setenv("ROCPROFILER_LOG", "1", 0);
setenv("ROCP_HSA_INTERCEPT", "1", 0);
setenv("ROCP_METRICS", _rocp_metrics.c_str(), 0);
setenv("HSA_TOOLS_REPORT_LOAD_FAILURE", "1", 0);
if(getenv("ROCM_PATH"))
{
setenv("OMNITRACE_ROCPROFILER_LIBRARY",
common::join('/', getenv("ROCM_PATH"),
"rocprofiler/lib/librocprofiler64.so")
.c_str(),
0);
}
#endif
#if OMNITRACE_USE_OMPT > 0
if(get_env("OMNITRACE_USE_OMPT", true))
{
projects/rocprofiler-systems/source/lib/omnitrace/CMakeLists.txt
+23 -2
Parādīt failu
@@ -32,6 +32,7 @@ target_link_libraries(
$<BUILD_INTERFACE:omnitrace::omnitrace-ptl>
$<BUILD_INTERFACE:omnitrace::omnitrace-hip>
$<BUILD_INTERFACE:omnitrace::omnitrace-roctracer>
$<BUILD_INTERFACE:omnitrace::omnitrace-rocprofiler>
$<BUILD_INTERFACE:omnitrace::omnitrace-rocm-smi>
$<BUILD_INTERFACE:$<IF:$<BOOL:${OMNITRACE_BUILD_LTO}>,omnitrace::omnitrace-lto,>>
$<BUILD_INTERFACE:$<IF:$<BOOL:${OMNITRACE_BUILD_STATIC_LIBGCC}>,omnitrace::omnitrace-static-libgcc,>>
@@ -98,6 +99,9 @@ set(library_headers
${CMAKE_CURRENT_LIST_DIR}/library/perfetto.hpp
${CMAKE_CURRENT_LIST_DIR}/library/process_sampler.hpp
${CMAKE_CURRENT_LIST_DIR}/library/ptl.hpp
${CMAKE_CURRENT_LIST_DIR}/library/rocm.hpp
${CMAKE_CURRENT_LIST_DIR}/library/rocprofiler.hpp
${CMAKE_CURRENT_LIST_DIR}/library/roctracer.hpp
${CMAKE_CURRENT_LIST_DIR}/library/runtime.hpp
${CMAKE_CURRENT_LIST_DIR}/library/sampling.hpp
${CMAKE_CURRENT_LIST_DIR}/library/state.hpp
@@ -113,8 +117,8 @@ set(library_headers
${CMAKE_CURRENT_LIST_DIR}/library/components/mpi_gotcha.hpp
${CMAKE_CURRENT_LIST_DIR}/library/components/omnitrace.hpp
${CMAKE_CURRENT_LIST_DIR}/library/components/rocm_smi.hpp
${CMAKE_CURRENT_LIST_DIR}/library/components/rocprofiler.hpp
${CMAKE_CURRENT_LIST_DIR}/library/components/roctracer.hpp
${CMAKE_CURRENT_LIST_DIR}/library/components/roctracer_callbacks.hpp
${CMAKE_CURRENT_LIST_DIR}/library/components/pthread_gotcha.hpp
${CMAKE_CURRENT_LIST_DIR}/library/components/pthread_create_gotcha.hpp
${CMAKE_CURRENT_LIST_DIR}/library/components/pthread_mutex_gotcha.hpp
@@ -122,11 +126,28 @@ set(library_headers
target_sources(omnitrace-object-library PRIVATE ${library_sources} ${library_headers})
if(OMNITRACE_USE_ROCTRACER OR OMNITRACE_USE_ROCPROFILER)
target_sources(
omnitrace-object-library
PRIVATE ${CMAKE_CURRENT_LIST_DIR}/library/rocprofiler.cpp
${CMAKE_CURRENT_LIST_DIR}/library/rocm.cpp
${CMAKE_CURRENT_LIST_DIR}/library/components/rocprofiler.cpp)
endif()
if(OMNITRACE_USE_ROCTRACER)
target_sources(
omnitrace-object-library
PRIVATE ${CMAKE_CURRENT_LIST_DIR}/library/components/roctracer.cpp
${CMAKE_CURRENT_LIST_DIR}/library/components/roctracer_callbacks.cpp)
${CMAKE_CURRENT_LIST_DIR}/library/roctracer.cpp)
endif()
if(OMNITRACE_USE_ROCPROFILER)
target_sources(
omnitrace-object-library
PRIVATE ${CMAKE_CURRENT_LIST_DIR}/library/rocprofiler.cpp
${CMAKE_CURRENT_LIST_DIR}/library/rocprofiler.hpp
${CMAKE_CURRENT_LIST_DIR}/library/rocprofiler/hsa_rsrc_factory.hpp
${CMAKE_CURRENT_LIST_DIR}/library/rocprofiler/hsa_rsrc_factory.cpp)
endif()
if(OMNITRACE_USE_ROCM_SMI)
projects/rocprofiler-systems/source/lib/omnitrace/library.cpp
+31 -2
Parādīt failu
@@ -29,6 +29,7 @@
#include "library/components/pthread_create_gotcha.hpp"
#include "library/components/pthread_gotcha.hpp"
#include "library/components/pthread_mutex_gotcha.hpp"
#include "library/components/rocprofiler.hpp"
#include "library/config.hpp"
#include "library/coverage.hpp"
#include "library/critical_trace.hpp"
@@ -38,6 +39,7 @@
#include "library/ompt.hpp"
#include "library/process_sampler.hpp"
#include "library/ptl.hpp"
#include "library/rocprofiler.hpp"
#include "library/sampling.hpp"
#include "library/thread_data.hpp"
#include "library/timemory.hpp"
@@ -91,6 +93,27 @@ ensure_finalization(bool _static_init = false)
tim::set_env("HSA_ENABLE_INTERRUPT", "0", 0);
#if defined(OMNITRACE_USE_ROCTRACER) && OMNITRACE_USE_ROCTRACER > 0
tim::set_env("HSA_TOOLS_LIB", "libomnitrace.so", 0);
#endif
#if defined(OMNITRACE_USE_ROCPROFILER) && OMNITRACE_USE_ROCPROFILER > 0
auto _default_rocm_path =
JOIN("", "/opt/rocm-", OMNITRACE_HIP_VERSION_MAJOR, '.',
OMNITRACE_HIP_VERSION_MINOR, '.', OMNITRACE_HIP_VERSION_PATCH);
auto _rocm_path = tim::get_env("OMNITRACE_ROCM_PATH",
tim::get_env("ROCM_PATH", _default_rocm_path));
auto _rocp_metrics = JOIN('/', _rocm_path, "rocprofiler/lib/metrics.xml");
tim::set_env("HSA_TOOLS_LIB", "libomnitrace.so", 0);
tim::set_env("ROCP_TOOL_LIB", "libomnitrace.so", 0);
tim::set_env("ROCPROFILER_LOG", "1", 0);
tim::set_env("ROCP_HSA_INTERCEPT", "1", 0);
tim::set_env("ROCP_METRICS", _rocp_metrics, 0);
tim::set_env("HSA_TOOLS_REPORT_LOAD_FAILURE", "1", 0);
if(getenv("ROCM_PATH"))
{
tim::set_env("OMNITRACE_ROCPROFILER_LIBRARY",
JOIN('/', tim::get_env<std::string>("ROCM_PATH"),
"rocprofiler/lib/librocprofiler64.so"),
0);
}
#endif
}
return scope::destructor{ []() { omnitrace_finalize_hidden(); } };
@@ -872,6 +895,13 @@ omnitrace_finalize_hidden(void)
tasking::join();
}
if(get_use_rocprofiler())
{
OMNITRACE_VERBOSE_F(1, "Shutting down rocprofiler...\n");
rocprofiler::post_process();
rocprofiler::rocm_cleanup();
}
if(dmp::rank() == 0) fprintf(stderr, "\n");
OMNITRACE_DEBUG_F("Stopping main bundle...\n");
@@ -1108,8 +1138,7 @@ omnitrace_finalize_hidden(void)
_push_count, "vs. popped:", _pop_count)
.c_str());
OMNITRACE_DEBUG_F("Disabling signal handling...\n");
tim::disable_signal_detection();
config::finalize();
OMNITRACE_VERBOSE_F(0, "Finalized\n");
}
projects/rocprofiler-systems/source/lib/omnitrace/library/common.hpp
+9 -2
Parādīt failu
@@ -28,6 +28,8 @@
#include <timemory/api.hpp>
#include <timemory/backends/dmp.hpp>
#include <timemory/backends/process.hpp>
#include <timemory/utility/demangle.hpp>
#include <timemory/utility/filepath.hpp>
#include <cassert>
#include <cstdint>
@@ -48,8 +50,13 @@ TIMEMORY_DEFINE_NS_API(api, rocm_smi)
namespace omnitrace
{
namespace api = tim::api; // NOLINT
namespace category = tim::category; // NOLINT
namespace api = ::tim::api; // NOLINT
namespace category = ::tim::category; // NOLINT
namespace filepath = ::tim::filepath; // NOLINT
using ::tim::demangle; // NOLINT
using ::tim::get_env; // NOLINT
using ::tim::try_demangle; // NOLINT
} // namespace omnitrace
// same sort of functionality as python's " ".join([...])
projects/rocprofiler-systems/source/lib/omnitrace/library/components/fwd.hpp
+13
Parādīt failu
@@ -39,6 +39,7 @@ TIMEMORY_DEFINE_NS_API(project, omnitrace)
TIMEMORY_DEFINE_NS_API(category, process_sampling)
TIMEMORY_DECLARE_COMPONENT(roctracer)
TIMEMORY_DECLARE_COMPONENT(rocprofiler)
/// \struct tim::trait::name
/// \brief provides a constexpr string in ::value
@@ -129,6 +130,7 @@ using sampling_gpu_temp = data_tracker<double, backtrace_gpu_temp>;
using sampling_gpu_power = data_tracker<double, backtrace_gpu_power>;
using sampling_gpu_memory = data_tracker<double, backtrace_gpu_memory>;
using roctracer = tim::component::roctracer;
using rocprofiler = tim::component::rocprofiler;
template <typename ApiT, typename StartFuncT = default_functor_t,
typename StopFuncT = default_functor_t>
@@ -140,6 +142,10 @@ struct functors;
TIMEMORY_DEFINE_CONCRETE_TRAIT(is_available, component::roctracer, false_type)
#endif
#if !defined(OMNITRACE_USE_ROCPROFILER)
TIMEMORY_DEFINE_CONCRETE_TRAIT(is_available, component::rocprofiler, false_type)
#endif
#if !defined(TIMEMORY_USE_LIBUNWIND)
TIMEMORY_DEFINE_CONCRETE_TRAIT(is_available, omnitrace::api::sampling, false_type)
TIMEMORY_DEFINE_CONCRETE_TRAIT(is_available, omnitrace::component::backtrace, false_type)
@@ -169,6 +175,9 @@ TIMEMORY_SET_COMPONENT_API(omnitrace::component::user_region, project::omnitrace
TIMEMORY_SET_COMPONENT_API(omnitrace::component::roctracer, project::omnitrace,
tpls::rocm, device::gpu, os::supports_linux,
category::external)
TIMEMORY_SET_COMPONENT_API(omnitrace::component::rocprofiler, project::omnitrace,
tpls::rocm, device::gpu, os::supports_linux,
category::external, category::hardware_counter)
TIMEMORY_SET_COMPONENT_API(omnitrace::component::sampling_wall_clock, project::omnitrace,
category::timing, os::supports_unix, category::sampling,
category::interrupt_sampling)
@@ -198,6 +207,8 @@ TIMEMORY_PROPERTY_SPECIALIZATION(omnitrace::component::user_region, OMNITRACE_US
"user_region", "omnitrace_user_region")
TIMEMORY_PROPERTY_SPECIALIZATION(omnitrace::component::roctracer, OMNITRACE_ROCTRACER,
"roctracer", "omnitrace_roctracer")
TIMEMORY_PROPERTY_SPECIALIZATION(omnitrace::component::rocprofiler, OMNITRACE_ROCPROFILER,
"rocprofiler", "omnitrace_rocprofiler")
TIMEMORY_PROPERTY_SPECIALIZATION(omnitrace::component::sampling_wall_clock,
OMNITRACE_SAMPLING_WALL_CLOCK, "sampling_wall_clock", "")
TIMEMORY_PROPERTY_SPECIALIZATION(omnitrace::component::sampling_cpu_clock,
@@ -226,6 +237,8 @@ TIMEMORY_METADATA_SPECIALIZATION(
"Used by OMPT")
TIMEMORY_METADATA_SPECIALIZATION(omnitrace::component::roctracer, "roctracer",
"High-precision ROCm API and kernel tracing", "")
TIMEMORY_METADATA_SPECIALIZATION(omnitrace::component::rocprofiler, "rocprofiler",
"ROCm kernel hardware counters", "")
TIMEMORY_METADATA_SPECIALIZATION(omnitrace::component::sampling_wall_clock,
"sampling_wall_clock", "Wall-clock timing",
"Derived from statistical sampling")
projects/rocprofiler-systems/source/lib/omnitrace/library/components/pthread_create_gotcha.cpp
+2 -1
Parādīt failu
@@ -154,7 +154,8 @@ pthread_create_gotcha::wrapper::operator()() const
if(!thread_bundle_data_t::instances().at(_tid))
{
thread_data<omnitrace_thread_bundle_t>::construct(
TIMEMORY_JOIN("", get_exe_name(), "/thread-", threading::get_id()),
TIMEMORY_JOIN('/', "omnitrace/process", process::get_id(), "thread",
threading::get_id()),
quirk::config<quirk::auto_start>{});
thread_bundle_data_t::instances().at(_tid)->start();
}
projects/rocprofiler-systems/source/lib/omnitrace/library/components/pthread_mutex_gotcha.cpp
+1 -1
Parādīt failu
@@ -29,8 +29,8 @@
#include "library/runtime.hpp"
#include "library/sampling.hpp"
#include "library/utility.hpp"
#include "timemory/backends/threading.hpp"
#include <timemory/backends/threading.hpp>
#include <timemory/utility/signals.hpp>
#include <timemory/utility/types.hpp>
projects/rocprofiler-systems/source/lib/omnitrace/library/components/rocprofiler.cpp
+344
Parādīt failu
@@ -0,0 +1,344 @@
// MIT License
//
// Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in all
// copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
// SOFTWARE.
#include "library/components/rocprofiler.hpp"
#include "library/common.hpp"
#include "library/components/pthread_create_gotcha.hpp"
#include "library/components/pthread_gotcha.hpp"
#include "library/config.hpp"
#include "library/debug.hpp"
#include "library/defines.hpp"
#include "library/dynamic_library.hpp"
#include "library/perfetto.hpp"
#include "library/redirect.hpp"
#include "library/rocprofiler.hpp"
#include "library/sampling.hpp"
#include "library/thread_data.hpp"
#include <timemory/storage/types.hpp>
#include <timemory/utility/types.hpp>
#include <timemory/variadic/functional.hpp>
#include <timemory/variadic/lightweight_tuple.hpp>
#include <rocprofiler.h>
#include <cstdint>
#include <string_view>
#include <type_traits>
TIMEMORY_STATISTICS_TYPE(component::rocm_data_tracker, component::rocm_feature_value)
TIMEMORY_DEFINE_CONCRETE_TRAIT(report_units, component::rocm_data_tracker, false_type)
namespace tim
{
namespace component
{
namespace
{
auto&
rocprofiler_activity_count()
{
static std::atomic<int64_t> _v{ 0 };
return _v;
}
} // namespace
omnitrace::unique_ptr_t<rocm_data_t>&
rocm_data(int64_t _tid)
{
using thread_data_t = omnitrace::thread_data<rocm_data_t, rocm_event>;
static auto& _v = thread_data_t::instances(thread_data_t::construct_on_init{});
return _v.at(_tid);
}
rocm_event::rocm_event(uint32_t _dev, uint32_t _thr, uint32_t _queue,
std::string _event_name, rocm_metric_type _begin,
rocm_metric_type _end, uint32_t _feature_count, void* _features_v)
: device_id{ _dev }
, thread_id{ _thr }
, queue_id{ _queue }
, entry{ _begin }
, exit{ _end }
, name(std::move(_event_name))
{
feature_values.reserve(_feature_count);
feature_names.reserve(_feature_count);
auto* _features = static_cast<rocprofiler_feature_t*>(_features_v);
for(uint32_t i = 0; i < _feature_count; ++i)
{
const rocprofiler_feature_t* p = &_features[i];
feature_names.emplace_back(p->name);
switch(p->data.kind)
{
// Output metrics results
case ROCPROFILER_DATA_KIND_UNINIT: break;
case ROCPROFILER_DATA_KIND_BYTES:
feature_values.emplace_back(
rocm_feature_value{ p->data.result_bytes.size });
break;
case ROCPROFILER_DATA_KIND_INT32:
feature_values.emplace_back(rocm_feature_value{ p->data.result_int32 });
break;
case ROCPROFILER_DATA_KIND_FLOAT:
feature_values.emplace_back(rocm_feature_value{ p->data.result_float });
break;
case ROCPROFILER_DATA_KIND_DOUBLE:
feature_values.emplace_back(rocm_feature_value{ p->data.result_double });
break;
case ROCPROFILER_DATA_KIND_INT64:
feature_values.emplace_back(rocm_feature_value{ p->data.result_int64 });
break;
}
}
}
std::string
rocm_event::as_string() const
{
std::stringstream _ss{};
_ss << name << ", device: " << device_id << ", queue: " << queue_id
<< ", thread: " << thread_id << ", entry: " << entry << ", exit = " << exit;
_ss.precision(3);
_ss << std::fixed;
for(size_t i = 0; i < feature_names.size(); ++i)
{
_ss << ", " << feature_names.at(i) << " = ";
auto _as_string = [&_ss](auto&& itr) { _ss << std::setw(4) << itr; };
std::visit(_as_string, feature_values.at(i));
}
return _ss.str();
}
void
rocprofiler::preinit()
{
rocprofiler_data::label() = "rocprofiler";
rocprofiler_data::description() = "ROCm hardware counters";
}
void
rocprofiler::start()
{
if(tracker_type::start() == 0) setup();
}
void
rocprofiler::stop()
{
if(tracker_type::stop() == 0) shutdown();
}
bool
rocprofiler::is_setup()
{
return omnitrace::rocprofiler::is_setup();
}
void
rocprofiler::add_setup(const std::string&, std::function<void()>&&)
{}
void
rocprofiler::add_shutdown(const std::string&, std::function<void()>&&)
{}
void
rocprofiler::remove_setup(const std::string&)
{}
void
rocprofiler::remove_shutdown(const std::string&)
{}
void
rocprofiler::setup()
{
OMNITRACE_VERBOSE_F(1, "rocprofiler is setup\n");
}
void
rocprofiler::shutdown()
{
omnitrace::rocprofiler::post_process();
omnitrace::rocprofiler::rocm_cleanup();
/*
using storage_type = typename rocprofiler_data::storage_type;
using bundle_t = rocprofiler_data;
using tag_t = api::omnitrace;
auto _data = omnitrace::rocprofiler::get_data();
auto _labels = omnitrace::rocprofiler::get_data_labels();
auto _info = omnitrace::rocprofiler::rocm_metrics();
int64_t _idx = 0;
auto _scope = tim::scope::get_default();
auto _get_metric_desc = [_info](std::string_view _v) {
for(auto itr : _info)
{
if(itr.symbol().find(_v) == 0 || itr.short_description().find(_v) == 0)
return std::make_pair(itr.short_description(), itr.long_description());
}
return std::make_pair(std::string{}, std::string{});
};
auto _debug = settings::debug();
settings::debug() = true;
struct hw_counters
{};
using rocm_counter = omnitrace::rocprofiler::rocm_counter;
struct perfetto_rocm_event
{
rocm_counter entry = {};
rocm_counter exit = {};
rocprofiler_value value = {};
bool operator<(const perfetto_rocm_event& _v) const
{
return (entry.at(0) == _v.entry.at(0)) ? exit.at(0) < _v.exit.at(0)
: entry.at(0) < _v.entry.at(0);
}
};
// contains the necessary info for export to perfetto
auto _perfetto_raw_data =
std::map<int64_t, std::map<int64_t, std::vector<perfetto_rocm_event>>>{};
// contains the time-stamp regions for the counter tracks
auto _perfetto_time_regions =
std::map<int64_t, std::map<int64_t, std::set<uint64_t>>>{};
// create a layout compatible for exporting to perfetto
for(const auto& itr : _labels)
{
auto _dev_id = itr.first;
auto _dev_name = JOIN("", '[', _dev_id, ']');
for(size_t i = 0; i < itr.second.size(); ++i)
{
auto _metric_name = itr.second.at(i);
auto _idx = perfetto_counter_track<hw_counters>::emplace(
_dev_id, JOIN(' ', "Device", _metric_name, _dev_name));
auto& _raw = _perfetto_raw_data[_dev_id][_idx];
auto& _reg = _perfetto_time_regions[_dev_id][_idx];
for(const auto& ditr : _data)
{
_raw.emplace_back(
perfetto_rocm_event{ ditr.entry, ditr.exit, ditr.data.at(i) });
}
std::sort(_raw.begin(), _raw.end());
for(auto ritr : _raw)
{
if(pthread_create_gotcha::is_valid_execution_time(0, ritr.entry.at(0)))
_reg.emplace(ritr.entry.at(0));
if(pthread_create_gotcha::is_valid_execution_time(0, ritr.exit.at(0)))
_reg.emplace(ritr.exit.at(0));
}
}
}
for(auto& ditr : _perfetto_time_regions)
for(auto& citr : ditr.second)
{
for(auto _ts = citr.second.begin(); _ts != citr.second.end(); ++_ts)
{
rocprofiler_value _v = {};
auto _curr = _ts;
auto _next = std::next(_ts);
if(_next == citr.second.end()) continue;
auto _min_ts = *_curr;
auto _max_ts = (_next == citr.second.end()) ? *_curr : *_next;
for(auto itr : _perfetto_raw_data[ditr.first][citr.first])
{
if(itr.entry[0] >= _min_ts && itr.exit[0] <= _max_ts)
{
using namespace tim::stl;
_v += itr.value;
}
}
auto _write_counter = [&](auto _v) {
if(_min_ts == _max_ts)
{
using value_type = std::remove_reference_t<
std::remove_cv_t<decay_t<decltype(_v)>>>;
_v = static_cast<value_type>(0);
}
TRACE_COUNTER(
"hardware_counter",
perfetto_counter_track<hw_counters>::at(ditr.first, citr.first),
_min_ts, _v);
};
std::visit(_write_counter, _v);
}
}
for(const auto& itr : _labels)
{
for(size_t i = 0; i < itr.second.size(); ++i)
{
auto _metric_name = itr.second.at(i);
auto _metric_desc = _get_metric_desc(_metric_name).second;
rocprofiler_data::label() = _metric_name;
if(!_metric_desc.empty())
rocprofiler_data::description() = JOIN(" - ", "rocprof", _metric_desc);
auto _dev_id = itr.first;
auto _label = JOIN('-', "rocprofiler", _metric_name, "device", _dev_id);
storage_type _storage{ standalone_storage{}, ++_idx, _label };
std::vector<bundle_t> _bundles = {};
_bundles.reserve(_data.size());
for(const auto& ditr : _data)
{
auto _hash = add_hash_id(ditr.name);
auto _v = ditr.data.at(i);
auto _obj = std::tie(_bundles.emplace_back(bundle_t{}));
invoke::reset<tag_t>(_obj);
invoke::push<tag_t>(_obj, _scope, _hash, &_storage, _dev_id);
invoke::start<tag_t>(_obj);
invoke::store<tag_t>(_obj, _v);
invoke::stop<tag_t>(_obj);
invoke::pop<tag_t>(_obj, &_storage, _dev_id);
}
_storage.write(_label);
}
}
settings::debug() = _debug;
*/
OMNITRACE_VERBOSE_F(1, "rocprofiler is shutdown\n");
}
scope::transient_destructor
rocprofiler::protect_flush_activity()
{
return scope::transient_destructor([]() { --rocprofiler_activity_count(); },
[]() { ++rocprofiler_activity_count(); });
}
} // namespace component
} // namespace tim
TIMEMORY_INSTANTIATE_EXTERN_COMPONENT(rocprofiler, false, void)
TIMEMORY_INSTANTIATE_EXTERN_COMPONENT(rocprofiler_data, true,
tim::component::rocprofiler_value)
projects/rocprofiler-systems/source/lib/omnitrace/library/components/rocprofiler.hpp
+236
Parādīt failu
@@ -0,0 +1,236 @@
// MIT License
//
// Copyright (c) 2022 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 "library/components/fwd.hpp"
#include "library/defines.hpp"
#include "library/thread_data.hpp"
#include <timemory/api.hpp>
#include <timemory/backends/hardware_counters.hpp>
#include <timemory/components/base.hpp>
#include <timemory/components/data_tracker/components.hpp>
#include <timemory/components/macros.hpp>
#include <timemory/enum.h>
#include <timemory/macros.hpp>
#include <timemory/macros/os.hpp>
#include <timemory/mpl/concepts.hpp>
#include <timemory/mpl/macros.hpp>
#include <timemory/mpl/type_traits.hpp>
#include <timemory/mpl/types.hpp>
#include <timemory/utility/transient_function.hpp>
#include <array>
#include <cstdint>
#include <string>
#include <string_view>
#include <variant>
#include <vector>
#if !defined(OMNITRACE_MAX_COUNTERS)
# define OMNITRACE_MAX_COUNTERS 25
#endif
#if !defined(OMNITRACE_ROCM_LOOK_AHEAD)
# define OMNITRACE_ROCM_LOOK_AHEAD 128
#endif
#if !defined(OMNITRACE_MAX_ROCM_QUEUES)
# define OMNITRACE_MAX_ROCM_QUEUES OMNITRACE_MAX_THREADS
#endif
namespace tim
{
namespace component
{
using rocm_metric_type = unsigned long long;
using rocm_info_entry = ::tim::hardware_counters::info;
using rocm_feature_value = std::variant<uint32_t, float, uint64_t, double>;
struct rocm_counter
{
std::array<rocm_metric_type, OMNITRACE_MAX_COUNTERS> counters;
};
struct rocm_event
{
using value_type = rocm_feature_value;
uint32_t device_id = 0;
uint32_t thread_id = 0;
uint32_t queue_id = 0;
rocm_metric_type entry = 0;
rocm_metric_type exit = 0;
std::string name = {};
std::vector<std::string_view> feature_names = {};
std::vector<rocm_feature_value> feature_values = {};
rocm_event() = default;
rocm_event(uint32_t _dev, uint32_t _thr, uint32_t _queue, std::string _event_name,
rocm_metric_type begin, rocm_metric_type end, uint32_t _feature_count,
void* _features);
std::string as_string() const;
friend std::ostream& operator<<(std::ostream& _os, const rocm_event& _v)
{
return (_os << _v.as_string());
}
friend bool operator<(const rocm_event& _lhs, const rocm_event& _rhs)
{
return std::tie(_lhs.device_id, _lhs.queue_id, _lhs.entry, _lhs.thread_id) <
std::tie(_rhs.device_id, _rhs.queue_id, _rhs.entry, _rhs.thread_id);
}
};
using rocm_data_t = std::vector<rocm_event>;
using rocm_data_tracker = data_tracker<rocm_feature_value, rocm_event>;
omnitrace::unique_ptr_t<rocm_data_t>&
rocm_data(int64_t _tid = threading::get_id());
using rocprofiler_value = typename rocm_event::value_type;
using rocprofiler_data = data_tracker<rocprofiler_value, rocprofiler>;
struct rocprofiler
: base<rocprofiler, void>
, private policy::instance_tracker<rocprofiler, false>
{
using value_type = void;
using base_type = base<rocprofiler, void>;
using tracker_type = policy::instance_tracker<rocprofiler, false>;
TIMEMORY_DEFAULT_OBJECT(rocprofiler)
static void preinit();
static void global_init() { setup(); }
static void global_finalize() { shutdown(); }
static bool is_setup();
static void setup();
static void shutdown();
static void add_setup(const std::string&, std::function<void()>&&);
static void add_shutdown(const std::string&, std::function<void()>&&);
static void remove_setup(const std::string&);
static void remove_shutdown(const std::string&);
void start();
void stop();
// this function protects rocprofiler_flush_activty from being called
// when omnitrace exits during a callback
[[nodiscard]] static scope::transient_destructor protect_flush_activity();
};
#if !defined(OMNITRACE_USE_ROCTRACER)
inline void
rocprofiler::setup()
{}
inline void
rocprofiler::shutdown()
{}
inline bool
rocprofiler::is_setup()
{
return false;
}
#endif
} // namespace component
namespace operation
{
template <>
struct set_storage<component::rocm_data_tracker>
{
using T = component::rocm_data_tracker;
static constexpr size_t max_threads = 4096;
using type = T;
using storage_array_t = std::array<storage<type>*, max_threads>;
friend struct get_storage<component::rocm_data_tracker>;
TIMEMORY_DEFAULT_OBJECT(set_storage)
auto operator()(storage<type>*, size_t) const {}
auto operator()(type&, size_t) const {}
auto operator()(storage<type>* _v) const { get().fill(_v); }
private:
static storage_array_t& get()
{
static storage_array_t _v = { nullptr };
return _v;
}
};
template <>
struct get_storage<component::rocm_data_tracker>
{
using type = component::rocm_data_tracker;
TIMEMORY_DEFAULT_OBJECT(get_storage)
auto operator()(const type&) const
{
return operation::set_storage<type>::get().at(0);
}
auto operator()() const
{
type _obj{};
return (*this)(_obj);
}
auto operator()(size_t _idx) const
{
return operation::set_storage<type>::get().at(_idx);
}
auto operator()(type&, size_t _idx) const { return (*this)(_idx); }
};
} // namespace operation
} // namespace tim
#if !defined(OMNITRACE_USE_ROCTRACER)
TIMEMORY_DEFINE_CONCRETE_TRAIT(is_available, component::rocprofiler_data, false_type)
#endif
TIMEMORY_SET_COMPONENT_API(component::rocprofiler_data, project::timemory,
category::timing, os::supports_unix)
TIMEMORY_DEFINE_CONCRETE_TRAIT(is_timing_category, component::rocprofiler_data,
false_type)
TIMEMORY_DEFINE_CONCRETE_TRAIT(uses_timing_units, component::rocprofiler_data, false_type)
TIMEMORY_DEFINE_CONCRETE_TRAIT(report_units, component::rocprofiler_data, false_type)
TIMEMORY_STATISTICS_TYPE(component::rocprofiler_data, component::rocprofiler_value)
#if !defined(OMNITRACE_EXTERN_COMPONENTS) || \
(defined(OMNITRACE_EXTERN_COMPONENTS) && OMNITRACE_EXTERN_COMPONENTS > 0)
# include <timemory/operations.hpp>
TIMEMORY_DECLARE_EXTERN_COMPONENT(rocprofiler, false, void)
TIMEMORY_DECLARE_EXTERN_COMPONENT(rocprofiler_data, true, double)
#endif
projects/rocprofiler-systems/source/lib/omnitrace/library/components/roctracer.cpp
+1 -1
Parādīt failu
@@ -22,12 +22,12 @@
#include "library/components/roctracer.hpp"
#include "library/components/pthread_gotcha.hpp"
#include "library/components/roctracer_callbacks.hpp"
#include "library/config.hpp"
#include "library/debug.hpp"
#include "library/defines.hpp"
#include "library/dynamic_library.hpp"
#include "library/redirect.hpp"
#include "library/roctracer.hpp"
#include "library/sampling.hpp"
#include "library/thread_data.hpp"
projects/rocprofiler-systems/source/lib/omnitrace/library/config.cpp
+64 -13
Parādīt failu
@@ -130,6 +130,14 @@ _settings_are_configured()
}
} // namespace
void
finalize()
{
OMNITRACE_DEBUG("[omnitrace_finalize] Disabling signal handling...\n");
tim::disable_signal_detection();
_settings_are_configured() = false;
}
bool
settings_are_configured()
{
@@ -142,7 +150,6 @@ configure_settings(bool _init)
{
volatile bool _v = _settings_are_configured();
if(_v) return;
_settings_are_configured() = true;
static bool _once = false;
if(_once) return;
@@ -161,6 +168,8 @@ configure_settings(bool _init)
tim::manager::add_metadata("OMNITRACE_VERSION_MAJOR", OMNITRACE_VERSION_MAJOR);
tim::manager::add_metadata("OMNITRACE_VERSION_MINOR", OMNITRACE_VERSION_MINOR);
tim::manager::add_metadata("OMNITRACE_VERSION_PATCH", OMNITRACE_VERSION_PATCH);
tim::manager::add_metadata("OMNITRACE_GIT_DESCRIBE", OMNITRACE_GIT_DESCRIBE);
tim::manager::add_metadata("OMNITRACE_GIT_REVISION", OMNITRACE_GIT_REVISION);
#if OMNITRACE_HIP_VERSION > 0
tim::manager::add_metadata("OMNITRACE_HIP_VERSION", OMNITRACE_HIP_VERSION_STRING);
@@ -206,8 +215,13 @@ configure_settings(bool _init)
!_config->get<bool>("OMNITRACE_USE_PERFETTO"), "backend",
"timemory");
OMNITRACE_CONFIG_SETTING(bool, "OMNITRACE_USE_ROCTRACER", "Enable ROCM tracing", true,
"backend", "roctracer", "rocm");
OMNITRACE_CONFIG_SETTING(bool, "OMNITRACE_USE_ROCTRACER",
"Enable ROCm API and kernel tracing", true, "backend",
"roctracer", "rocm");
OMNITRACE_CONFIG_SETTING(bool, "OMNITRACE_USE_ROCPROFILER",
"Enable ROCm hardware counters", true, "backend",
"rocprofiler", "rocm");
OMNITRACE_CONFIG_SETTING(
bool, "OMNITRACE_USE_ROCM_SMI",
@@ -343,6 +357,13 @@ configure_settings(bool _init)
OMNITRACE_CONFIG_SETTING(std::string, "OMNITRACE_ROCTRACER_HSA_API_TYPES",
"HSA API type to collect", "", "roctracer", "rocm");
OMNITRACE_CONFIG_SETTING(
std::string, "OMNITRACE_ROCM_EVENTS",
"ROCm hardware counters. Use ':device=N' syntax to specify collection on device "
"number N, e.g. ':device=0'. If no device specification is provided, the event "
"is collected on every available device",
"", "rocprofiler", "rocm", "hardware_counters");
OMNITRACE_CONFIG_SETTING(bool, "OMNITRACE_CRITICAL_TRACE_DEBUG",
"Enable debugging for critical trace", _omnitrace_debug,
"debugging", "critical_trace");
@@ -553,6 +574,8 @@ configure_settings(bool _init)
configure_signal_handler();
configure_disabled_settings();
_settings_are_configured() = true;
OMNITRACE_CONDITIONAL_BASIC_PRINT(get_verbose_env() > 0, "configuration complete\n");
}
@@ -582,6 +605,7 @@ configure_mode_settings()
_set("OMNITRACE_USE_TIMEMORY", false);
_set("OMNITRACE_USE_ROCM_SMI", false);
_set("OMNITRACE_USE_ROCTRACER", false);
_set("OMNITRACE_USE_ROCPROFILER", false);
_set("OMNITRACE_USE_KOKKOSP", false);
_set("OMNITRACE_USE_OMPT", false);
_set("OMNITRACE_USE_SAMPLING", false);
@@ -599,8 +623,9 @@ configure_mode_settings()
{
OMNITRACE_VERBOSE_F(
1, "No HIP devices were found: disabling roctracer and rocm_smi...\n");
get_use_roctracer() = false;
get_use_rocm_smi() = false;
_set("OMNITRACE_USE_ROCTRACER", false);
_set("OMNITRACE_USE_ROCPROFILER", false);
_set("OMNITRACE_USE_ROCM_SMI", false);
}
get_instrumentation_interval() = std::max<size_t>(get_instrumentation_interval(), 1);
@@ -624,6 +649,7 @@ configure_mode_settings()
_set("OMNITRACE_USE_TIMEMORY", false);
_set("OMNITRACE_USE_ROCM_SMI", false);
_set("OMNITRACE_USE_ROCTRACER", false);
_set("OMNITRACE_USE_ROCPROFILER", false);
_set("OMNITRACE_USE_KOKKOSP", false);
_set("OMNITRACE_USE_OMPT", false);
_set("OMNITRACE_USE_SAMPLING", false);
@@ -723,6 +749,7 @@ configure_disabled_settings()
_handle_use_option("OMNITRACE_USE_OMPT", "ompt");
_handle_use_option("OMNITRACE_USE_ROCM_SMI", "rocm_smi");
_handle_use_option("OMNITRACE_USE_ROCTRACER", "roctracer");
_handle_use_option("OMNITRACE_USE_ROCPROFILER", "rocprofiler");
_handle_use_option("OMNITRACE_CRITICAL_TRACE", "critical_trace");
#if !defined(OMNITRACE_USE_ROCTRACER) || OMNITRACE_USE_ROCTRACER == 0
@@ -731,6 +758,12 @@ configure_disabled_settings()
_config->find(itr)->second->set_hidden(true);
#endif
#if !defined(OMNITRACE_USE_ROCPROFILER) || OMNITRACE_USE_ROCPROFILER == 0
_config->find("OMNITRACE_USE_ROCPROFILER")->second->set_hidden(true);
for(const auto& itr : _config->disable_category("rocprofiler"))
_config->find(itr)->second->set_hidden(true);
#endif
#if !defined(OMNITRACE_USE_ROCM_SMI) || OMNITRACE_USE_ROCM_SMI == 0
_config->find("OMNITRACE_USE_ROCM_SMI")->second->set_hidden(true);
for(const auto& itr : _config->disable_category("rocm_smi"))
@@ -1064,7 +1097,8 @@ is_binary_rewrite()
bool
get_debug_env()
{
return tim::get_env<bool>("OMNITRACE_DEBUG", false);
return (settings_are_configured()) ? get_debug()
: tim::get_env<bool>("OMNITRACE_DEBUG", false);
}
bool
@@ -1106,7 +1140,8 @@ get_debug_sampling()
int
get_verbose_env()
{
return tim::get_env<int>("OMNITRACE_VERBOSE", 0);
return (settings_are_configured()) ? get_verbose()
: tim::get_env<int>("OMNITRACE_VERBOSE", 0);
}
int
@@ -1130,27 +1165,36 @@ get_use_timemory()
return static_cast<tim::tsettings<bool>&>(*_v->second).get();
}
bool&
bool
get_use_roctracer()
{
#if defined(OMNITRACE_USE_ROCTRACER) && OMNITRACE_USE_ROCTRACER > 0
static auto _v = get_config()->find("OMNITRACE_USE_ROCTRACER");
return static_cast<tim::tsettings<bool>&>(*_v->second).get();
#else
static auto _v = false;
return _v;
return false;
#endif
}
bool&
bool
get_use_rocprofiler()
{
#if defined(OMNITRACE_USE_ROCPROFILER) && OMNITRACE_USE_ROCPROFILER > 0
static auto _v = get_config()->find("OMNITRACE_USE_ROCPROFILER");
return static_cast<tim::tsettings<bool>&>(*_v->second).get();
#else
return false;
#endif
}
bool
get_use_rocm_smi()
{
#if defined(OMNITRACE_USE_ROCM_SMI) && OMNITRACE_USE_ROCM_SMI > 0
static auto _v = get_config()->find("OMNITRACE_USE_ROCM_SMI");
return static_cast<tim::tsettings<bool>&>(*_v->second).get();
#else
static auto _v = false;
return _v;
return false;
#endif
}
@@ -1442,6 +1486,13 @@ get_trace_thread_locks()
return static_cast<tim::tsettings<bool>&>(*_v->second).get();
}
std::string
get_rocm_events()
{
static auto _v = get_config()->find("OMNITRACE_ROCM_EVENTS");
return static_cast<tim::tsettings<std::string>&>(*_v->second).get();
}
bool
get_trace_thread_rwlocks()
{
projects/rocprofiler-systems/source/lib/omnitrace/library/config.hpp
+11 -2
Parādīt failu
@@ -57,6 +57,9 @@ configure_signal_handler();
void
configure_disabled_settings();
void
finalize();
void
handle_deprecated_setting(const std::string& _old, const std::string& _new,
int _verbose = 0);
@@ -162,10 +165,13 @@ get_use_perfetto() OMNITRACE_HOT;
bool&
get_use_timemory() OMNITRACE_HOT;
bool&
bool
get_use_roctracer() OMNITRACE_HOT;
bool&
bool
get_use_rocprofiler() OMNITRACE_HOT;
bool
get_use_rocm_smi() OMNITRACE_HOT;
bool&
@@ -276,6 +282,9 @@ get_trace_thread_locks();
bool
get_trace_thread_rwlocks();
std::string
get_rocm_events();
} // namespace config
//
projects/rocprofiler-systems/source/lib/omnitrace/library/defines.hpp.in
+11 -4
Parādīt failu
@@ -27,6 +27,8 @@
#define OMNITRACE_VERSION_MAJOR @PROJECT_VERSION_MAJOR@
#define OMNITRACE_VERSION_MINOR @PROJECT_VERSION_MINOR@
#define OMNITRACE_VERSION_PATCH @PROJECT_VERSION_PATCH@
#define OMNITRACE_GIT_DESCRIBE "@OMNITRACE_GIT_DESCRIBE@"
#define OMNITRACE_GIT_REVISION "@OMNITRACE_GIT_REVISION@"
#define OMNITRACE_HIP_VERSION_STRING "@OMNITRACE_HIP_VERSION@"
#define OMNITRACE_HIP_VERSION_MAJOR @OMNITRACE_HIP_VERSION_MAJOR@
@@ -40,20 +42,25 @@
#endif
// clang-format on
#define OMNITRACE_VERSION \
((10000 * OMNITRACE_VERSION_MAJOR) + (100 * OMNITRACE_VERSION_MINOR) + \
OMNITRACE_VERSION_PATCH)
#define OMNITRACE_HIP_VERSION \
((10000 * OMNITRACE_HIP_VERSION_MAJOR) + (100 * OMNITRACE_HIP_VERSION_MINOR) + \
OMNITRACE_HIP_VERSION_PATCH)
#define TIMEMORY_USER_COMPONENT_ENUM \
OMNITRACE_COMPONENT_idx, OMNITRACE_USER_REGION_idx, OMNITRACE_ROCTRACER_idx, \
OMNITRACE_SAMPLING_WALL_CLOCK_idx, OMNITRACE_SAMPLING_CPU_CLOCK_idx, \
OMNITRACE_SAMPLING_PERCENT_idx, OMNITRACE_SAMPLING_GPU_POWER_idx, \
OMNITRACE_SAMPLING_GPU_TEMP_idx, OMNITRACE_SAMPLING_GPU_BUSY_idx, \
OMNITRACE_SAMPLING_GPU_MEMORY_USAGE_idx,
OMNITRACE_ROCPROFILER_idx, OMNITRACE_SAMPLING_WALL_CLOCK_idx, \
OMNITRACE_SAMPLING_CPU_CLOCK_idx, OMNITRACE_SAMPLING_PERCENT_idx, \
OMNITRACE_SAMPLING_GPU_POWER_idx, OMNITRACE_SAMPLING_GPU_TEMP_idx, \
OMNITRACE_SAMPLING_GPU_BUSY_idx, OMNITRACE_SAMPLING_GPU_MEMORY_USAGE_idx,
#define OMNITRACE_COMPONENT OMNITRACE_COMPONENT_idx
#define OMNITRACE_USER_REGION OMNITRACE_USER_REGION_idx
#define OMNITRACE_ROCTRACER OMNITRACE_ROCTRACER_idx
#define OMNITRACE_ROCPROFILER OMNITRACE_ROCPROFILER_idx
#define OMNITRACE_SAMPLING_WALL_CLOCK OMNITRACE_SAMPLING_WALL_CLOCK_idx
#define OMNITRACE_SAMPLING_CPU_CLOCK OMNITRACE_SAMPLING_CPU_CLOCK_idx
#define OMNITRACE_SAMPLING_PERCENT OMNITRACE_SAMPLING_PERCENT_idx
projects/rocprofiler-systems/source/lib/omnitrace/library/dynamic_library.hpp
+21
Parādīt failu
@@ -26,6 +26,7 @@
#include <dlfcn.h>
#include <string>
#include <unistd.h>
namespace omnitrace
{
@@ -45,6 +46,26 @@ struct dynamic_library
bool open();
int close() const;
template <typename RetT, typename... Args>
RetT invoke(std::string_view _name, RetT (*&_func)(Args...), Args... _args)
{
if(!handle) open();
if(handle)
{
*(void**) (&_func) = dlsym(handle, _name.data());
if(_func)
{
return (*_func)(_args...);
}
else
{
fprintf(stderr, "[omnitrace][pid=%i]> %s :: %s\n", getpid(), _name.data(),
dlerror());
}
}
return RetT{};
}
std::string envname = {};
std::string filename = {};
int flags = 0;
projects/rocprofiler-systems/source/lib/omnitrace/library/gpu.cpp
+19 -3
Parādīt failu
@@ -24,25 +24,41 @@
#if defined(OMNITRACE_USE_ROCM_SMI) && OMNITRACE_USE_ROCM_SMI > 0
# include "library/components/rocm_smi.hpp"
#elif defined(OMNITRACE_USE_HIP) && OMNITRACE_USE_HIP > 0
#elif !defined(OMNITRACE_USE_ROCM_SMI)
# define OMNITRACE_USE_ROCM_SMI 0
#endif
#if defined(OMNITRACE_USE_HIP) && OMNITRACE_USE_HIP > 0
# if !defined(TIMEMORY_USE_HIP)
# define TIMEMORY_USE_HIP 1
# endif
# include <timemory/components/hip/backends.hpp>
#elif !defined(OMNITRACE_USE_HIP)
# define OMNITRACE_USE_HIP 0
#endif
namespace omnitrace
{
namespace gpu
{
int
hip_device_count()
{
#if OMNITRACE_USE_HIP > 0
return ::tim::hip::device_count();
#else
return 0;
#endif
}
int
device_count()
{
#if defined(OMNITRACE_USE_ROCM_SMI) && OMNITRACE_USE_ROCM_SMI > 0
#if OMNITRACE_USE_ROCM_SMI > 0
// store as static since calls after rsmi_shutdown will return zero
static auto _v = rocm_smi::device_count();
return _v;
#elif defined(OMNITRACE_USE_HIP) && OMNITRACE_USE_HIP > 0
#elif OMNITRACE_USE_HIP > 0
return ::tim::hip::device_count();
#else
return 0;
projects/rocprofiler-systems/source/lib/omnitrace/library/gpu.hpp
+4 -1
Parādīt failu
@@ -28,5 +28,8 @@ namespace gpu
{
int
device_count();
}
int
hip_device_count();
} // namespace gpu
} // namespace omnitrace
projects/rocprofiler-systems/source/lib/omnitrace/library/ompt.cpp
+2 -2
Parādīt failu
@@ -68,8 +68,8 @@ setup()
comp::user_ompt_bundle::reset();
tim::auto_lock_t lk{ tim::type_mutex<ompt_handle_t>() };
comp::user_ompt_bundle::configure<omnitrace::component::user_region>();
f_bundle =
std::make_unique<ompt_bundle_t>("ompt", quirk::config<quirk::auto_start>{});
f_bundle = std::make_unique<ompt_bundle_t>("omnitrace/ompt",
quirk::config<quirk::auto_start>{});
}
void
projects/rocprofiler-systems/source/lib/omnitrace/library/perfetto.hpp
+3 -1
Parādīt failu
@@ -173,7 +173,8 @@ struct perfetto_counter_track
_missing.emplace_back(std::make_tuple(*itr, itr->c_str(), false));
}
}
auto& _name = _name_data.emplace_back(std::make_unique<std::string>(_v));
auto _index = _track_data.size();
auto& _name = _name_data.emplace_back(std::make_unique<std::string>(_v));
const char* _unit_name = (_units && strlen(_units) > 0) ? _units : nullptr;
_track_data.emplace_back(perfetto::CounterTrack{ _name->c_str() }
.set_unit_name(_unit_name)
@@ -217,6 +218,7 @@ struct perfetto_counter_track
}
}
}
return _index;
}
static auto& at(size_t _idx, size_t _n) { return get_data().second.at(_idx).at(_n); }
projects/rocprofiler-systems/source/lib/omnitrace/library/rocm.cpp
+264
Parādīt failu
@@ -0,0 +1,264 @@
// MIT License
//
// Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in all
// copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
// SOFTWARE.
#include "library/rocm.hpp"
#include "library.hpp"
#include "library/components/rocm_smi.hpp"
#include "library/components/rocprofiler.hpp"
#include "library/components/roctracer.hpp"
#include "library/config.hpp"
#include "library/critical_trace.hpp"
#include "library/debug.hpp"
#include "library/rocprofiler.hpp"
#include "library/rocprofiler/hsa_rsrc_factory.hpp"
#include "library/roctracer.hpp"
#include "library/runtime.hpp"
#include "library/sampling.hpp"
#include "library/thread_data.hpp"
#include "library/tracing.hpp"
#include <timemory/backends/cpu.hpp>
#include <timemory/backends/threading.hpp>
#include <timemory/utility/types.hpp>
#include <atomic>
#include <chrono>
#include <cstdint>
#include <mutex>
#include <tuple>
#define HIP_PROF_HIP_API_STRING 1
#include <roctracer_ext.h>
#include <roctracer_hcc.h>
#include <roctracer_hip.h>
#define AMD_INTERNAL_BUILD 1
#include <roctracer_hsa.h>
#if __has_include(<hip/amd_detail/hip_prof_str.h>) || (defined(OMNITRACE_USE_HIP) && OMNITRACE_USE_HIP > 0)
# include <hip/amd_detail/hip_prof_str.h>
# define OMNITRACE_HIP_API_ARGS 1
#else
# define OMNITRACE_HIP_API_ARGS 0
#endif
#include <rocprofiler.h>
using namespace omnitrace;
namespace omnitrace
{
namespace rocm
{
std::mutex rocm_mutex = {};
bool is_loaded = false;
} // namespace rocm
} // namespace omnitrace
// HSA-runtime tool on-load method
extern "C"
{
#if defined(OMNITRACE_USE_ROCPROFILER) && OMNITRACE_USE_ROCPROFILER > 0
void OnUnloadTool()
{
OMNITRACE_BASIC_VERBOSE(2, "Inside %s\n", __FUNCTION__);
rocm::lock_t _lk{ rocm::rocm_mutex, std::defer_lock };
if(!_lk.owns_lock()) _lk.lock();
if(!rocm::is_loaded) return;
rocm::is_loaded = false;
_lk.unlock();
// stop_top_level_timer_if_necessary();
// Final resources cleanup
omnitrace::rocprofiler::rocm_cleanup();
}
void OnLoadToolProp(rocprofiler_settings_t* settings)
{
OMNITRACE_BASIC_VERBOSE(2, "Inside %s\n", __FUNCTION__);
rocm::lock_t _lk{ rocm::rocm_mutex, std::defer_lock };
if(!_lk.owns_lock()) _lk.lock();
if(rocm::is_loaded) return;
rocm::is_loaded = true;
_lk.unlock();
// Enable timestamping
settings->timestamp_on = 1u;
// Initialize profiling
omnitrace::rocprofiler::rocm_initialize();
HsaRsrcFactory::Instance().PrintGpuAgents("ROCm");
}
#endif
bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
const char* const* failed_tool_names)
{
OMNITRACE_BASIC_VERBOSE(2, "Inside %s\n", __FUNCTION__);
if(!tim::get_env("OMNITRACE_INIT_TOOLING", true)) return true;
if(!tim::settings::enabled()) return true;
roctracer_is_init() = true;
pthread_gotcha::push_enable_sampling_on_child_threads(false);
OMNITRACE_BASIC_VERBOSE_F(1, "\n");
tim::consume_parameters(table, runtime_version, failed_tool_count,
failed_tool_names);
if(!config::settings_are_configured() && get_state() < State::Active)
omnitrace_init_tooling_hidden();
OMNITRACE_SCOPED_THREAD_STATE(ThreadState::Internal);
static auto _setup = [=]() {
try
{
OMNITRACE_VERBOSE(1, "[OnLoad] setting up HSA...\n");
// const char* output_prefix = getenv("ROCP_OUTPUT_DIR");
const char* output_prefix = nullptr;
bool trace_hsa_api = get_trace_hsa_api();
// Enable HSA API callbacks/activity
if(trace_hsa_api)
{
std::vector<std::string> hsa_api_vec =
tim::delimit(get_trace_hsa_api_types());
// initialize HSA tracing
roctracer_set_properties(ACTIVITY_DOMAIN_HSA_API, (void*) table);
OMNITRACE_VERBOSE(1, " HSA-trace(");
if(!hsa_api_vec.empty())
{
for(const auto& itr : hsa_api_vec)
{
uint32_t cid = HSA_API_ID_NUMBER;
const char* api = itr.c_str();
ROCTRACER_CALL(roctracer_op_code(ACTIVITY_DOMAIN_HSA_API, api,
&cid, nullptr));
ROCTRACER_CALL(roctracer_enable_op_callback(
ACTIVITY_DOMAIN_HSA_API, cid, hsa_api_callback, nullptr));
OMNITRACE_VERBOSE(1, " %s", api);
}
}
else
{
ROCTRACER_CALL(roctracer_enable_domain_callback(
ACTIVITY_DOMAIN_HSA_API, hsa_api_callback, nullptr));
}
OMNITRACE_VERBOSE(1, " )\n");
}
bool trace_hsa_activity = get_trace_hsa_activity();
// Enable HSA GPU activity
if(trace_hsa_activity)
{
// initialize HSA tracing
::roctracer::hsa_ops_properties_t ops_properties{
table,
reinterpret_cast<activity_async_callback_t>(
hsa_activity_callback),
nullptr, output_prefix
};
roctracer_set_properties(ACTIVITY_DOMAIN_HSA_OPS, &ops_properties);
OMNITRACE_VERBOSE(1, " HSA-activity-trace()\n");
ROCTRACER_CALL(roctracer_enable_op_activity(ACTIVITY_DOMAIN_HSA_OPS,
HSA_OP_ID_COPY));
}
} catch(std::exception& _e)
{
OMNITRACE_BASIC_PRINT("Exception was thrown in HSA setup: %s\n",
_e.what());
}
};
static auto _shutdown = []() {
OMNITRACE_DEBUG_F("roctracer_disable_domain_callback\n");
ROCTRACER_CALL(roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HSA_API));
OMNITRACE_DEBUG_F("roctracer_disable_op_activity\n");
ROCTRACER_CALL(
roctracer_disable_op_activity(ACTIVITY_DOMAIN_HSA_OPS, HSA_OP_ID_COPY));
};
(void) omnitrace::get_clock_skew();
comp::roctracer::add_setup("hsa", _setup);
comp::roctracer::add_shutdown("hsa", _shutdown);
rocm_smi::set_state(State::Active);
comp::roctracer::setup();
#if defined(OMNITRACE_USE_ROCPROFILER) && OMNITRACE_USE_ROCPROFILER > 0
bool _force_rocprofiler_init =
tim::get_env("OMNITRACE_FORCE_ROCPROFILE_INIT", false, false);
#else
bool _force_rocprofiler_init = false;
#endif
bool _success = true;
bool _is_empty =
(config::settings_are_configured() && config::get_rocm_events().empty());
if(_force_rocprofiler_init || (get_use_rocprofiler() && !_is_empty))
{
auto _rocprof =
dynamic_library{ "OMNITRACE_ROCPROFILER_LIBRARY", "librocprofiler64.so",
(RTLD_LAZY | RTLD_GLOBAL), true };
on_load_t _rocprof_load = nullptr;
_success = _rocprof.invoke("OnLoad", _rocprof_load, table, runtime_version,
failed_tool_count, failed_tool_names);
OMNITRACE_CONDITIONAL_PRINT_F(!_success,
"Warning! Invoking rocprofiler's OnLoad "
"failed! OMNITRACE_ROCPROFILER_LIBRARY=%s\n",
_rocprof.filename.c_str());
OMNITRACE_CI_THROW(!_success,
"Warning! Invoking rocprofiler's OnLoad "
"failed! OMNITRACE_ROCPROFILER_LIBRARY=%s\n",
_rocprof.filename.c_str());
}
pthread_gotcha::pop_enable_sampling_on_child_threads();
return _success;
}
// HSA-runtime on-unload method
void OnUnload()
{
OMNITRACE_BASIC_VERBOSE(2, "Inside %s\n", __FUNCTION__);
rocm_smi::set_state(State::Finalized);
comp::roctracer::shutdown();
comp::rocprofiler::shutdown();
omnitrace_finalize_hidden();
}
}
projects/rocprofiler-systems/source/lib/omnitrace/library/rocm.hpp
+58
Parādīt failu
@@ -0,0 +1,58 @@
// MIT License
//
// Copyright (c) 2022 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 "library/defines.hpp"
#if defined(OMNITRACE_USE_ROCPROFILER) && OMNITRACE_USE_ROCPROFILER > 0
# include <rocprofiler.h>
#endif
#include <cstdint>
#include <mutex>
namespace omnitrace
{
namespace rocm
{
using lock_t = std::unique_lock<std::mutex>;
extern std::mutex rocm_mutex;
extern bool is_loaded;
} // namespace rocm
} // namespace omnitrace
extern "C"
{
struct HsaApiTable;
using on_load_t = bool (*)(HsaApiTable*, uint64_t, uint64_t, const char* const*);
bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
const char* const* failed_tool_names) OMNITRACE_PUBLIC_API;
void OnUnload() OMNITRACE_PUBLIC_API;
#if defined(OMNITRACE_USE_ROCPROFILER) && OMNITRACE_USE_ROCPROFILER > 0
void OnLoadToolProp(rocprofiler_settings_t* settings) OMNITRACE_PUBLIC_API;
void OnUnloadTool() OMNITRACE_PUBLIC_API;
#endif
}
projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler.cpp
+791
Parādīt failu
@@ -0,0 +1,791 @@
// MIT License
//
// Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in all
// copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
// SOFTWARE.
#include "library/rocprofiler.hpp"
#include "library/common.hpp"
#include "library/config.hpp"
#include "library/debug.hpp"
#include "library/gpu.hpp"
#include "library/perfetto.hpp"
#include "library/rocm.hpp"
#include "library/rocprofiler/hsa_rsrc_factory.hpp"
#include <timemory/backends/hardware_counters.hpp>
#include <timemory/manager.hpp>
#include <timemory/mpl/concepts.hpp>
#include <timemory/storage/types.hpp>
#include <timemory/utility/types.hpp>
#include <rocprofiler.h>
#include <atomic>
#include <cstdlib>
#include <dlfcn.h>
#include <hsa.h>
#include <iostream>
#include <mutex>
#include <sstream>
#include <string.h>
#include <string_view>
#include <type_traits>
#include <unistd.h>
#include <vector>
namespace omnitrace
{
namespace rocprofiler
{
namespace
{
auto&
get_event_names()
{
static auto _v = std::map<uint32_t, std::vector<rocprofiler_feature_t>>{};
return _v;
}
} // namespace
// Error handler
void
fatal(const std::string& msg)
{
OMNITRACE_PRINT_F("\n");
OMNITRACE_PRINT_F("%s\n", msg.c_str());
abort();
}
// Check returned HSA API status
const char*
rocm_error_string(hsa_status_t _status)
{
const char* _err_string = nullptr;
if(_status != HSA_STATUS_SUCCESS) rocprofiler_error_string(&_err_string);
return _err_string;
}
// Check returned HSA API status
bool
rocm_check_status(hsa_status_t _status, const std::set<hsa_status_t>& _nonfatal = {})
{
if(_status != HSA_STATUS_SUCCESS)
{
if(_nonfatal.count(_status) == 0)
fatal(JOIN(" :: ", "ERROR", rocm_error_string(_status)));
OMNITRACE_PRINT_F("Warning! %s\n", rocm_error_string(_status));
return false;
}
return true;
}
// Context stored entry type
struct context_entry_t
{
bool valid;
hsa_agent_t agent;
rocprofiler_group_t group;
rocprofiler_callback_data_t data;
};
// Context callback arg
struct callbacks_arg_t
{
rocprofiler_pool_t** pools;
};
// Handler callback arg
struct handler_arg_t
{
rocprofiler_feature_t* features;
unsigned feature_count;
};
bool&
is_setup()
{
static bool _v = false;
return _v;
}
std::map<uint32_t, std::vector<std::string_view>>
get_data_labels()
{
auto _v = std::map<uint32_t, std::vector<std::string_view>>{};
for(const auto& itr : get_event_names())
{
_v[itr.first] = {};
for(auto vitr : itr.second)
_v[itr.first].emplace_back(std::string_view{ vitr.name });
}
return _v;
}
// Dump stored context entry
void
rocm_dump_context_entry(context_entry_t* entry, rocprofiler_feature_t* features,
unsigned feature_count)
{
// static rocm_metric_type last_timestamp = get_last_timestamp_ns();
volatile std::atomic<bool>* valid =
reinterpret_cast<std::atomic<bool>*>(&entry->valid);
while(valid->load() == false)
sched_yield();
const rocprofiler_dispatch_record_t* record = entry->data.record;
if(!record) return; // there is nothing to do here.
auto _queue_id = entry->data.queue_id;
auto _thread_id = entry->data.thread_id;
auto _dev_id = HsaRsrcFactory::Instance().GetAgentInfo(entry->agent)->dev_index;
auto _kernel_name = std::string{ entry->data.kernel_name };
auto _pos = _kernel_name.find_last_of(')');
if(_pos != std::string::npos) _kernel_name = _kernel_name.substr(0, _pos + 1);
rocprofiler_group_t& group = entry->group;
if(group.context == nullptr)
{
fatal("context is nullptr\n");
}
if(feature_count > 0)
{
rocm_check_status(rocprofiler_group_get_data(&group));
rocm_check_status(rocprofiler_get_metrics(group.context));
}
auto _evt = comp::rocm_event{ _dev_id, _thread_id, _queue_id, _kernel_name,
record->begin, record->end, feature_count, features };
comp::rocm_data()->emplace_back(_evt);
}
// Profiling completion handler
// Dump and delete the context entry
// Return true if the context was dumped successfully
bool
rocm_context_handler(const rocprofiler_pool_entry_t* entry, void* arg)
{
// Context entry
context_entry_t* ctx_entry = reinterpret_cast<context_entry_t*>(entry->payload);
handler_arg_t* handler_arg = reinterpret_cast<handler_arg_t*>(arg);
rocm::lock_t _lk{ rocm::rocm_mutex, std::defer_lock };
if(!_lk.owns_lock()) _lk.lock();
rocm_dump_context_entry(ctx_entry, handler_arg->features, handler_arg->feature_count);
return true;
}
// Kernel disoatch callback
hsa_status_t
rocm_dispatch_callback(const rocprofiler_callback_data_t* callback_data, void* arg,
rocprofiler_group_t* group)
{
// Passed tool data
hsa_agent_t agent = callback_data->agent;
// Open profiling context
const unsigned gpu_id = HsaRsrcFactory::Instance().GetAgentInfo(agent)->dev_index;
callbacks_arg_t* callbacks_arg = reinterpret_cast<callbacks_arg_t*>(arg);
rocprofiler_pool_t* pool = callbacks_arg->pools[gpu_id];
rocprofiler_pool_entry_t pool_entry{};
rocm_check_status(rocprofiler_pool_fetch(pool, &pool_entry));
// Profiling context entry
rocprofiler_t* context = pool_entry.context;
context_entry_t* entry = reinterpret_cast<context_entry_t*>(pool_entry.payload);
// Get group[0]
rocm_check_status(rocprofiler_get_group(context, 0, group));
// Fill profiling context entry
entry->agent = agent;
entry->group = *group;
entry->data = *callback_data;
entry->data.kernel_name = strdup(callback_data->kernel_name);
reinterpret_cast<std::atomic<bool>*>(&entry->valid)->store(true);
return HSA_STATUS_SUCCESS;
}
unsigned
metrics_input(unsigned _device, rocprofiler_feature_t** ret)
{
// OMNITRACE_THROW("%s\n", __FUNCTION__);
// Profiling feature objects
auto _events = tim::delimit(config::get_rocm_events(), ", ;\t\n");
std::vector<std::string> _features = {};
auto _this_device = JOIN("", ":device=", _device);
for(auto itr : _events)
{
OMNITRACE_VERBOSE_F(3, "Processing feature '%s' for device %u...\n", itr.c_str(),
_device);
auto _pos = itr.find(":device=");
if(_pos != std::string::npos)
{
if(itr.find(_this_device) != std::string::npos)
{
_features.emplace_back(itr.substr(0, _pos));
}
}
else
{
_features.emplace_back(itr);
}
}
const unsigned feature_count = _features.size();
rocprofiler_feature_t* features = new rocprofiler_feature_t[feature_count];
memset(features, 0, feature_count * sizeof(rocprofiler_feature_t));
// PMC events
for(unsigned i = 0; i < feature_count; ++i)
{
OMNITRACE_VERBOSE_F(3, "Adding feature '%s' for device %u...\n",
_features.at(i).c_str(), _device);
features[i].kind = ROCPROFILER_FEATURE_KIND_METRIC;
features[i].name = strdup(_features.at(i).c_str());
features[i].parameters = nullptr;
features[i].parameter_count = 0;
}
*ret = features;
return feature_count;
}
struct info_data
{
const AgentInfo* agent = nullptr;
std::vector<comp::rocm_info_entry>* data = nullptr;
};
hsa_status_t
info_data_callback(const rocprofiler_info_data_t info, void* arg)
{
using qualifier_t = tim::hardware_counters::qualifier;
using qualifier_vec_t = std::vector<qualifier_t>;
auto* _arg = static_cast<info_data*>(arg);
const auto* _agent = _arg->agent;
auto* _data = _arg->data;
switch(info.kind)
{
case ROCPROFILER_INFO_KIND_METRIC:
{
auto _device_qualifier_sym = JOIN("", ":device=", _agent->dev_index);
auto _device_qualifier = tim::hardware_counters::qualifier{
true, static_cast<int>(_agent->dev_index), _device_qualifier_sym,
JOIN(" ", "Device", _agent->dev_index)
};
auto _long_desc = std::string{ info.metric.description };
auto _units = std::string{};
auto _pysym = std::string{};
if(info.metric.expr != nullptr)
{
auto _sym = JOIN("", info.metric.name, _device_qualifier_sym);
auto _short_desc = JOIN("", "Derived counter: ", info.metric.expr);
_data->emplace_back(comp::rocm_info_entry(
true, tim::hardware_counters::api::rocm, _data->size(), 0, _sym,
_pysym, _short_desc, _long_desc, _units,
qualifier_vec_t{ _device_qualifier }));
}
else
{
if(info.metric.instances == 1)
{
auto _sym = JOIN("", info.metric.name, _device_qualifier_sym);
auto _short_desc =
JOIN("", info.metric.name, " on device ", _agent->dev_index);
_data->emplace_back(comp::rocm_info_entry(
true, tim::hardware_counters::api::rocm, _data->size(), 0, _sym,
_pysym, _short_desc, _long_desc, _units,
qualifier_vec_t{ _device_qualifier }));
}
else
{
for(uint32_t i = 0; i < info.metric.instances; ++i)
{
auto _instance_qualifier_sym = JOIN("", '[', i, ']');
auto _instance_qualifier =
tim::hardware_counters::qualifier{ true, static_cast<int>(i),
_instance_qualifier_sym,
JOIN(" ", "Instance", i) };
auto _sym = JOIN("", info.metric.name, _instance_qualifier_sym,
_device_qualifier_sym);
auto _short_desc = JOIN("", info.metric.name, " instance ", i,
" on device ", _agent->dev_index);
_data->emplace_back(comp::rocm_info_entry(
true, tim::hardware_counters::api::rocm, _data->size(), 0,
_sym, _pysym, _short_desc, _long_desc, _units,
qualifier_vec_t{ _device_qualifier, _instance_qualifier }));
}
}
}
break;
}
default: printf("wrong info kind %u\n", info.kind); return HSA_STATUS_ERROR;
}
return HSA_STATUS_SUCCESS;
}
std::vector<comp::rocm_info_entry>
rocm_metrics()
{
std::vector<comp::rocm_info_entry> _data = {};
try
{
(void) HsaRsrcFactory::Instance();
} catch(std::runtime_error& _e)
{
OMNITRACE_VERBOSE_F(0, "%s\n", _e.what());
return _data;
}
// Available GPU agents
const unsigned gpu_count = HsaRsrcFactory::Instance().GetCountOfGpuAgents();
std::vector<AgentInfo*> _gpu_agents(gpu_count, nullptr);
for(unsigned i = 0; i < gpu_count; ++i)
{
const AgentInfo* _agent = _gpu_agents[i];
const AgentInfo** _agent_p = &_agent;
HsaRsrcFactory::Instance().GetGpuAgentInfo(i, _agent_p);
auto _v = info_data{ _agent, &_data };
if(!rocm_check_status(
rocprofiler_iterate_info(&_agent->dev_id, ROCPROFILER_INFO_KIND_METRIC,
info_data_callback, reinterpret_cast<void*>(&_v)),
{ HSA_STATUS_ERROR_NOT_INITIALIZED }))
return _data;
}
auto _settings = tim::settings::shared_instance();
if(_settings)
{
auto ritr = _settings->find("OMNITRACE_ROCM_EVENTS");
if(ritr != _settings->end())
{
auto _rocm_events = ritr->second;
if(_rocm_events->get_choices().empty())
{
std::vector<std::string> _choices = {};
_choices.reserve(_data.size());
for(auto itr : _data)
{
if(!itr.symbol().empty()) _choices.emplace_back(itr.symbol());
}
_rocm_events->set_choices(_choices);
}
}
}
return _data;
}
void
rocm_initialize()
{
// Available GPU agents
const unsigned gpu_count = HsaRsrcFactory::Instance().GetCountOfGpuAgents();
(void) rocm_metrics();
// Adding dispatch observer
callbacks_arg_t* callbacks_arg = new callbacks_arg_t{};
callbacks_arg->pools = new rocprofiler_pool_t*[gpu_count];
for(unsigned gpu_id = 0; gpu_id < gpu_count; gpu_id++)
{
// Getting profiling features
rocprofiler_feature_t* features = nullptr;
unsigned feature_count = metrics_input(gpu_id, &features);
if(features)
{
get_event_names()[gpu_id].clear();
get_event_names()[gpu_id].reserve(feature_count);
for(unsigned i = 0; i < feature_count; ++i)
get_event_names().at(gpu_id).emplace_back(features[i]);
}
// Handler arg
handler_arg_t* handler_arg = new handler_arg_t{};
handler_arg->features = features;
handler_arg->feature_count = feature_count;
// Context properties
rocprofiler_pool_properties_t properties{};
properties.num_entries = 100;
properties.payload_bytes = sizeof(context_entry_t);
properties.handler = rocm_context_handler;
properties.handler_arg = handler_arg;
// Getting GPU device info
const AgentInfo* agent_info = nullptr;
if(HsaRsrcFactory::Instance().GetGpuAgentInfo(gpu_id, &agent_info) == false)
{
fprintf(stderr, "GetGpuAgentInfo failed\n");
abort();
}
// Open profiling pool
rocprofiler_pool_t* pool = nullptr;
uint32_t mode = 0; // ROCPROFILER_MODE_SINGLEGROUP
rocm_check_status(rocprofiler_pool_open(agent_info->dev_id, features,
feature_count, &pool, mode, &properties));
callbacks_arg->pools[gpu_id] = pool;
}
rocprofiler_queue_callbacks_t callbacks_ptrs{};
callbacks_ptrs.dispatch = rocm_dispatch_callback;
int err = rocprofiler_set_queue_callbacks(callbacks_ptrs, callbacks_arg);
OMNITRACE_VERBOSE_F(3, "err=%d, rocprofiler_set_queue_callbacks\n", err);
is_setup() = true;
}
void
rocm_cleanup()
{
// Unregister dispatch callback
rocm_check_status(rocprofiler_remove_queue_callbacks());
// close profiling pool
// rocm_check_status(rocprofiler_pool_flush(pool));
// rocm_check_status(rocprofiler_pool_close(pool));
}
namespace
{
using rocm_event = comp::rocm_event;
using rocm_data_t = comp::rocm_data_t;
using rocm_metric_type = comp::rocm_metric_type;
using rocm_feature_value = comp::rocm_feature_value;
using rocm_data_tracker = comp::rocm_data_tracker;
void
post_process_perfetto()
{
using counter_track = perfetto_counter_track<rocm_event>;
static bool _once = false;
if(_once) return;
auto _data = rocm_data_t{};
auto _device_data = std::map<uint32_t, std::vector<rocm_event*>>{};
auto _device_fields = std::map<uint32_t, std::vector<std::string_view>>{};
auto _device_range = std::map<uint32_t, std::set<rocm_metric_type>>{};
for(size_t i = 0; i < OMNITRACE_MAX_THREADS; ++i)
{
auto& _v = comp::rocm_data(i);
if(_v)
{
_data.reserve(_data.size() + _v->size());
for(auto& itr : *_v)
_data.emplace_back(itr);
}
}
if(_data.empty()) return;
_once = true;
std::sort(_data.begin(), _data.end());
auto _get_events = [](std::vector<rocm_event*>& _inp, rocm_metric_type _ts) {
auto _v = std::vector<rocm_event*>{};
for(const auto& itr : _inp)
{
if(_ts >= itr->entry && _ts <= itr->exit) _v.emplace_back(itr);
}
return _v;
};
{
auto _device_time = std::map<uint32_t, std::set<rocm_metric_type>>{};
for(auto& itr : _data)
{
_device_data[itr.device_id].emplace_back(&itr);
_device_time[itr.device_id].emplace(itr.entry);
_device_time[itr.device_id].emplace(itr.exit);
auto _dev_id = itr.device_id;
if(get_use_perfetto() && !counter_track::exists(_dev_id))
{
auto addendum = [&](auto&& _v) {
return JOIN(" ", "Device", _v, JOIN("", '[', _dev_id, ']'));
};
for(auto nitr : itr.feature_names)
counter_track::emplace(_dev_id, addendum(nitr));
}
}
for(auto& ditr : _device_time)
{
for(auto itr = ditr.second.begin(); itr != ditr.second.end(); ++itr)
{
auto _next = std::next(itr);
if(_next == ditr.second.end()) continue;
_device_range[ditr.first].emplace(((*_next / 2) + (*itr / 2)));
}
}
}
for(auto& ditr : _device_range)
{
auto _dev_id = ditr.first;
auto _values = std::vector<rocm_feature_value>{};
for(const auto& itr : ditr.second)
{
auto _v = _get_events(_device_data[_dev_id], itr);
uint64_t _ts = itr;
for(auto* vitr : _v)
{
size_t _n = vitr->feature_values.size();
if(_values.empty())
{
_values.reserve(_n);
for(size_t i = 0; i < _n; ++i)
{
_values.emplace_back(vitr->feature_values.at(i));
}
}
else
{
for(size_t i = 0; i < _n; ++i)
{
auto _plus = [](auto& _lhs, auto&& _rhs) { _lhs += _rhs; };
std::visit(_plus, _values.at(i), vitr->feature_values.at(i));
}
}
}
for(size_t i = 0; i < _values.size(); ++i)
{
auto _trace_counter = [_dev_id, i, _ts](auto&& _v) {
TRACE_COUNTER("hardware_counter", counter_track::at(_dev_id, i), _ts,
_v);
};
std::visit(_trace_counter, _values.at(i));
}
}
}
}
void
post_process_timemory()
{
static bool _once = false;
if(_once) return;
auto _data = rocm_data_t{};
auto _device_data = std::map<uint32_t, std::vector<rocm_event*>>{};
auto _device_fields = std::map<uint32_t, std::vector<std::string_view>>{};
auto _device_range = std::map<uint32_t, std::set<rocm_metric_type>>{};
for(size_t i = 0; i < OMNITRACE_MAX_THREADS; ++i)
{
auto& _v = comp::rocm_data(i);
if(_v)
{
_data.reserve(_data.size() + _v->size());
for(auto& itr : *_v)
_data.emplace_back(itr);
}
}
if(_data.empty()) return;
_once = true;
std::sort(_data.begin(), _data.end());
for(auto& itr : _data)
{
_device_data[itr.device_id].emplace_back(&itr);
}
using storage_type = typename rocm_data_tracker::storage_type;
using bundle_type = tim::lightweight_tuple<rocm_data_tracker>;
auto _info = rocm_metrics();
static auto _get_description = [&_info](std::string_view _v) {
for(auto& itr : _info)
{
if(itr.symbol().find(_v) == 0 || itr.short_description().find(_v) == 0)
{
return itr.long_description();
}
}
return std::string{};
};
struct local_event
{
rocm_event* parent = nullptr;
std::vector<local_event> children = {};
TIMEMORY_DEFAULT_OBJECT(local_event)
explicit local_event(rocm_event* _v)
: parent{ _v }
{}
bool operator()(rocm_event* _v)
{
OMNITRACE_CI_THROW(!parent, "Error! '%s' has nullptr", __PRETTY_FUNCTION__);
if(_v->device_id != parent->device_id) return false;
if(_v->entry > parent->entry && _v->exit <= parent->exit)
{
for(auto& itr : children)
{
if(itr(_v)) return true;
}
children.emplace_back(_v);
std::sort(children.begin(), children.end());
return true;
}
return false;
}
bool operator<(const local_event& _v) const
{
OMNITRACE_CI_THROW(!parent, "Error! '%s' has nullptr", __PRETTY_FUNCTION__);
OMNITRACE_CI_THROW(!_v.parent, "Error! '%s' passed nullptr",
__PRETTY_FUNCTION__);
return *parent < *_v.parent;
}
void operator()(int64_t _index, scope::config _scope) const
{
OMNITRACE_CI_THROW(!parent, "Error! '%s' has nullptr", __PRETTY_FUNCTION__);
bundle_type _bundle{ parent->name, _scope };
_bundle.push(parent->queue_id)
.start()
.store(parent->feature_values.at(_index));
for(const auto& itr : children)
itr(_index, _scope);
_bundle.stop().pop(parent->queue_id);
}
};
struct local_storage
{
int64_t index = 0;
std::string metric_name = {};
std::string metric_description = {};
std::unique_ptr<storage_type> storage = {};
local_storage(uint32_t _devid, size_t _idx, std::string_view _name)
: index{ static_cast<int64_t>(_idx) }
, metric_name{ _name }
, metric_description{ _get_description(metric_name) }
{
auto _metric_name = std::string{ _name };
_metric_name = std::regex_replace(
_metric_name, std::regex{ "(.*)\\[([0-9]+)\\]" }, "$1_$2");
storage = std::make_unique<storage_type>(
tim::standalone_storage{}, index,
JOIN('-', "rocprof", "device", _devid, _metric_name));
}
void operator()(const local_event& _event, scope::config _scope) const
{
operation::set_storage<rocm_data_tracker>{}(storage.get());
_event(index, _scope);
}
void write() const
{
rocm_data_tracker::label() = metric_name;
rocm_data_tracker::description() = metric_description;
storage->write();
}
};
auto _local_data = std::map<uint32_t, std::vector<local_event>>{};
auto _scope = scope::get_default();
for(auto& ditr : _device_data)
{
auto _storage = std::vector<local_storage>{};
for(auto& itr : ditr.second)
{
auto _n = itr->feature_names.size();
if(_n > _storage.size())
{
_storage.reserve(_n);
for(size_t i = _storage.size(); i < _n; ++i)
_storage.emplace_back(ditr.first, i, itr->feature_names.at(i));
}
}
auto& _local = _local_data[ditr.first];
for(auto& itr : ditr.second)
{
for(auto& litr : _local)
{
if(litr(itr))
{
goto _bypass_insert;
}
}
_local.emplace_back(itr);
_bypass_insert:;
}
for(auto& sitr : _storage)
{
for(auto& itr : _local_data[ditr.first])
sitr(itr, _scope);
}
for(auto& itr : _storage)
itr.write();
}
tim::trait::runtime_enabled<omnitrace::rocprofiler::rocm_data_tracker>::set(false);
}
} // namespace
void
post_process()
{
if(get_use_perfetto()) post_process_perfetto();
if(get_use_timemory())
{
auto _manager = tim::manager::master_instance();
if(_manager)
{
_manager->add_cleanup("rocprofiler", &post_process_timemory);
}
else
{
post_process_timemory();
}
}
}
} // namespace rocprofiler
} // namespace omnitrace
projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler.hpp
+88
Parādīt failu
@@ -0,0 +1,88 @@
// MIT License
//
// Copyright (c) 2022 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 "library/components/rocprofiler.hpp"
#include "library/defines.hpp"
#include "library/timemory.hpp"
#include <timemory/backends/hardware_counters.hpp>
#include <timemory/macros.hpp>
#include <timemory/mpl/concepts.hpp>
#include <timemory/mpl/macros.hpp>
#include <array>
#include <atomic>
#include <cstring>
#include <dlfcn.h>
#include <iostream>
#include <list>
#include <map>
#include <string>
#include <string_view>
#include <tuple>
#include <unistd.h>
#include <utility>
#include <variant>
#include <vector>
namespace omnitrace
{
namespace rocprofiler
{
std::map<uint32_t, std::vector<std::string_view>>
get_data_labels();
void
rocm_initialize();
void
rocm_cleanup();
bool&
is_setup();
void
post_process();
std::vector<comp::rocm_info_entry>
rocm_metrics();
#if !defined(OMNITRACE_USE_ROCPROFILER) || OMNITRACE_USE_ROCPROFILER == 0
inline void
post_process()
{}
inline void
rocm_cleanup()
{}
inline std::vector<comp::rocm_info_entry>
rocm_metrics()
{
return std::vector<comp::rocm_info_entry>{};
}
#endif
} // namespace rocprofiler
} // namespace omnitrace
projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler/hsa_rsrc_factory.cpp
+976
Parādīt failu
@@ -0,0 +1,976 @@
/******************************************************************************
Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*******************************************************************************/
#include "library/rocprofiler/hsa_rsrc_factory.hpp"
#include "library/debug.hpp"
#include <timemory/manager.hpp>
#include <rocprofiler.h>
#define PUBLIC_API __attribute__((visibility("default")))
#define CONSTRUCTOR_API __attribute__((constructor))
#define DESTRUCTOR_API __attribute__((destructor))
#include <atomic>
#include <cassert>
#include <dlfcn.h>
#include <fcntl.h>
#include <fstream>
#include <hsa.h>
#include <hsa_ext_amd.h>
#include <hsa_ext_finalize.h>
#include <iostream>
#include <stdint.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <string>
#include <sys/mman.h>
#include <sys/stat.h>
#include <sys/types.h>
#include <vector>
// Callback function to get available in the system agents
hsa_status_t
HsaRsrcFactory::GetHsaAgentsCallback(hsa_agent_t agent, void* data)
{
hsa_status_t status = HSA_STATUS_ERROR;
HsaRsrcFactory* hsa_rsrc = reinterpret_cast<HsaRsrcFactory*>(data);
const AgentInfo* agent_info = hsa_rsrc->AddAgentInfo(agent);
if(agent_info != nullptr) status = HSA_STATUS_SUCCESS;
return status;
}
// This function checks to see if the provided
// pool has the HSA_AMD_SEGMENT_GLOBAL property. If the kern_arg flag is true,
// the function adds an additional requirement that the pool have the
// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT property. If kern_arg is false,
// pools must NOT have this property.
// Upon finding a pool that meets these conditions, HSA_STATUS_INFO_BREAK is
// returned. HSA_STATUS_SUCCESS is returned if no errors were encountered, but
// no pool was found meeting the requirements. If an error is encountered, we
// return that error.
static hsa_status_t
FindGlobalPool(hsa_amd_memory_pool_t pool, void* data, bool kern_arg)
{
hsa_status_t err;
hsa_amd_segment_t segment;
uint32_t flag;
if(nullptr == data)
{
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
err = HsaRsrcFactory::HsaApi()->hsa_amd_memory_pool_get_info(
pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment);
CHECK_STATUS("hsa_amd_memory_pool_get_info", err);
if(HSA_AMD_SEGMENT_GLOBAL != segment)
{
return HSA_STATUS_SUCCESS;
}
err = HsaRsrcFactory::HsaApi()->hsa_amd_memory_pool_get_info(
pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag);
CHECK_STATUS("hsa_amd_memory_pool_get_info", err);
uint32_t karg_st = flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT;
if((karg_st == 0 && kern_arg) || (karg_st != 0 && !kern_arg))
{
return HSA_STATUS_SUCCESS;
}
*(reinterpret_cast<hsa_amd_memory_pool_t*>(data)) = pool;
return HSA_STATUS_INFO_BREAK;
}
// This is the call-back function for hsa_amd_agent_iterate_memory_pools() that
// finds a pool with the properties of HSA_AMD_SEGMENT_GLOBAL and that is NOT
// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT
hsa_status_t
FindStandardPool(hsa_amd_memory_pool_t pool, void* data)
{
return FindGlobalPool(pool, data, false);
}
// This is the call-back function for hsa_amd_agent_iterate_memory_pools() that
// finds a pool with the properties of HSA_AMD_SEGMENT_GLOBAL and that IS
// HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT
hsa_status_t
FindKernArgPool(hsa_amd_memory_pool_t pool, void* data)
{
return FindGlobalPool(pool, data, true);
}
// Constructor of the class
HsaRsrcFactory::HsaRsrcFactory(bool initialize_hsa)
: initialize_hsa_(initialize_hsa)
{
hsa_status_t status;
cpu_pool_ = nullptr;
kern_arg_pool_ = nullptr;
InitHsaApiTable(nullptr);
// Initialize the Hsa Runtime
if(initialize_hsa_)
{
status = hsa_api_.hsa_init();
CHECK_STATUS("Error in hsa_init", status);
}
// Discover the set of Gpu devices available on the platform
status = hsa_api_.hsa_iterate_agents(GetHsaAgentsCallback, this);
CHECK_STATUS("Error Calling hsa_iterate_agents", status);
if(cpu_pool_ == nullptr)
CHECK_STATUS("CPU memory pool is not found", HSA_STATUS_ERROR);
if(kern_arg_pool_ == nullptr)
CHECK_STATUS("Kern-arg memory pool is not found", HSA_STATUS_ERROR);
// Get AqlProfile API table
aqlprofile_api_ = { nullptr };
#ifdef ROCP_LD_AQLPROFILE
status = LoadAqlProfileLib(&aqlprofile_api_);
#else
status = hsa_api_.hsa_system_get_major_extension_table(
HSA_EXTENSION_AMD_AQLPROFILE, hsa_ven_amd_aqlprofile_VERSION_MAJOR,
sizeof(aqlprofile_api_), &aqlprofile_api_);
#endif
CHECK_STATUS("aqlprofile API table load failed", status);
// Get Loader API table
loader_api_ = { nullptr };
status = hsa_api_.hsa_system_get_major_extension_table(
HSA_EXTENSION_AMD_LOADER, 1, sizeof(loader_api_), &loader_api_);
CHECK_STATUS("loader API table query failed", status);
// Instantiate HSA timer
timer_ = new HsaTimer(&hsa_api_);
CHECK_STATUS("HSA timer allocation failed",
(timer_ == nullptr) ? HSA_STATUS_ERROR : HSA_STATUS_SUCCESS);
// System timeout
timeout_ = (timeout_ns_ == HsaTimer::TIMESTAMP_MAX)
? timeout_ns_
: timer_->ns_to_sysclock(timeout_ns_);
}
// Destructor of the class
HsaRsrcFactory::~HsaRsrcFactory()
{
delete timer_;
for(const auto* p : cpu_list_)
delete p;
for(const auto* p : gpu_list_)
delete p;
if(initialize_hsa_)
{
hsa_status_t status = hsa_api_.hsa_shut_down();
try
{
CHECK_STATUS("Error in hsa_shut_down", status);
} catch(std::runtime_error& _e)
{
fflush(stderr);
fprintf(stderr, "%s\n", _e.what());
fflush(stderr);
abort();
}
}
}
void
HsaRsrcFactory::InitHsaApiTable(HsaApiTable* table)
{
std::lock_guard<mutex_t> lck(mutex_);
if(hsa_api_.hsa_init == nullptr)
{
if(table != nullptr)
{
hsa_api_.hsa_init = table->core_->hsa_init_fn;
hsa_api_.hsa_shut_down = table->core_->hsa_shut_down_fn;
hsa_api_.hsa_agent_get_info = table->core_->hsa_agent_get_info_fn;
hsa_api_.hsa_iterate_agents = table->core_->hsa_iterate_agents_fn;
hsa_api_.hsa_queue_create = table->core_->hsa_queue_create_fn;
hsa_api_.hsa_queue_destroy = table->core_->hsa_queue_destroy_fn;
hsa_api_.hsa_queue_load_write_index_relaxed =
table->core_->hsa_queue_load_write_index_relaxed_fn;
hsa_api_.hsa_queue_store_write_index_relaxed =
table->core_->hsa_queue_store_write_index_relaxed_fn;
hsa_api_.hsa_queue_load_read_index_relaxed =
table->core_->hsa_queue_load_read_index_relaxed_fn;
hsa_api_.hsa_signal_create = table->core_->hsa_signal_create_fn;
hsa_api_.hsa_signal_destroy = table->core_->hsa_signal_destroy_fn;
hsa_api_.hsa_signal_load_relaxed = table->core_->hsa_signal_load_relaxed_fn;
hsa_api_.hsa_signal_store_relaxed = table->core_->hsa_signal_store_relaxed_fn;
hsa_api_.hsa_signal_wait_scacquire =
table->core_->hsa_signal_wait_scacquire_fn;
hsa_api_.hsa_signal_store_screlease =
table->core_->hsa_signal_store_screlease_fn;
hsa_api_.hsa_code_object_reader_create_from_file =
table->core_->hsa_code_object_reader_create_from_file_fn;
hsa_api_.hsa_executable_create_alt =
table->core_->hsa_executable_create_alt_fn;
hsa_api_.hsa_executable_load_agent_code_object =
table->core_->hsa_executable_load_agent_code_object_fn;
hsa_api_.hsa_executable_freeze = table->core_->hsa_executable_freeze_fn;
hsa_api_.hsa_executable_get_symbol =
table->core_->hsa_executable_get_symbol_fn;
hsa_api_.hsa_executable_symbol_get_info =
table->core_->hsa_executable_symbol_get_info_fn;
hsa_api_.hsa_executable_iterate_symbols =
table->core_->hsa_executable_iterate_symbols_fn;
hsa_api_.hsa_system_get_info = table->core_->hsa_system_get_info_fn;
hsa_api_.hsa_system_get_major_extension_table =
table->core_->hsa_system_get_major_extension_table_fn;
hsa_api_.hsa_amd_agent_iterate_memory_pools =
table->amd_ext_->hsa_amd_agent_iterate_memory_pools_fn;
hsa_api_.hsa_amd_memory_pool_get_info =
table->amd_ext_->hsa_amd_memory_pool_get_info_fn;
hsa_api_.hsa_amd_memory_pool_allocate =
table->amd_ext_->hsa_amd_memory_pool_allocate_fn;
hsa_api_.hsa_amd_agents_allow_access =
table->amd_ext_->hsa_amd_agents_allow_access_fn;
hsa_api_.hsa_amd_memory_async_copy =
table->amd_ext_->hsa_amd_memory_async_copy_fn;
hsa_api_.hsa_amd_signal_async_handler =
table->amd_ext_->hsa_amd_signal_async_handler_fn;
hsa_api_.hsa_amd_profiling_set_profiler_enabled =
table->amd_ext_->hsa_amd_profiling_set_profiler_enabled_fn;
hsa_api_.hsa_amd_profiling_get_async_copy_time =
table->amd_ext_->hsa_amd_profiling_get_async_copy_time_fn;
hsa_api_.hsa_amd_profiling_get_dispatch_time =
table->amd_ext_->hsa_amd_profiling_get_dispatch_time_fn;
}
else
{
hsa_api_.hsa_init = hsa_init;
hsa_api_.hsa_shut_down = hsa_shut_down;
hsa_api_.hsa_agent_get_info = hsa_agent_get_info;
hsa_api_.hsa_iterate_agents = hsa_iterate_agents;
hsa_api_.hsa_queue_create = hsa_queue_create;
hsa_api_.hsa_queue_destroy = hsa_queue_destroy;
hsa_api_.hsa_queue_load_write_index_relaxed =
hsa_queue_load_write_index_relaxed;
hsa_api_.hsa_queue_store_write_index_relaxed =
hsa_queue_store_write_index_relaxed;
hsa_api_.hsa_queue_load_read_index_relaxed =
hsa_queue_load_read_index_relaxed;
hsa_api_.hsa_signal_create = hsa_signal_create;
hsa_api_.hsa_signal_destroy = hsa_signal_destroy;
hsa_api_.hsa_signal_load_relaxed = hsa_signal_load_relaxed;
hsa_api_.hsa_signal_store_relaxed = hsa_signal_store_relaxed;
hsa_api_.hsa_signal_wait_scacquire = hsa_signal_wait_scacquire;
hsa_api_.hsa_signal_store_screlease = hsa_signal_store_screlease;
hsa_api_.hsa_code_object_reader_create_from_file =
hsa_code_object_reader_create_from_file;
hsa_api_.hsa_executable_create_alt = hsa_executable_create_alt;
hsa_api_.hsa_executable_load_agent_code_object =
hsa_executable_load_agent_code_object;
hsa_api_.hsa_executable_freeze = hsa_executable_freeze;
hsa_api_.hsa_executable_get_symbol = hsa_executable_get_symbol;
hsa_api_.hsa_executable_symbol_get_info = hsa_executable_symbol_get_info;
hsa_api_.hsa_executable_iterate_symbols = hsa_executable_iterate_symbols;
hsa_api_.hsa_system_get_info = hsa_system_get_info;
hsa_api_.hsa_system_get_major_extension_table =
hsa_system_get_major_extension_table;
hsa_api_.hsa_amd_agent_iterate_memory_pools =
hsa_amd_agent_iterate_memory_pools;
hsa_api_.hsa_amd_memory_pool_get_info = hsa_amd_memory_pool_get_info;
hsa_api_.hsa_amd_memory_pool_allocate = hsa_amd_memory_pool_allocate;
hsa_api_.hsa_amd_agents_allow_access = hsa_amd_agents_allow_access;
hsa_api_.hsa_amd_memory_async_copy = hsa_amd_memory_async_copy;
hsa_api_.hsa_amd_signal_async_handler = hsa_amd_signal_async_handler;
hsa_api_.hsa_amd_profiling_set_profiler_enabled =
hsa_amd_profiling_set_profiler_enabled;
hsa_api_.hsa_amd_profiling_get_async_copy_time =
hsa_amd_profiling_get_async_copy_time;
hsa_api_.hsa_amd_profiling_get_dispatch_time =
hsa_amd_profiling_get_dispatch_time;
}
}
}
hsa_status_t
HsaRsrcFactory::LoadAqlProfileLib(aqlprofile_pfn_t* api)
{
void* handle = dlopen(kAqlProfileLib, RTLD_NOW);
if(handle == nullptr)
{
fprintf(stderr, "Loading '%s' failed, %s\n", kAqlProfileLib, dlerror());
return HSA_STATUS_ERROR;
}
dlerror(); /* Clear any existing error */
api->hsa_ven_amd_aqlprofile_error_string =
(decltype(::hsa_ven_amd_aqlprofile_error_string)*) dlsym(
handle, "hsa_ven_amd_aqlprofile_error_string");
api->hsa_ven_amd_aqlprofile_validate_event =
(decltype(::hsa_ven_amd_aqlprofile_validate_event)*) dlsym(
handle, "hsa_ven_amd_aqlprofile_validate_event");
api->hsa_ven_amd_aqlprofile_start = (decltype(::hsa_ven_amd_aqlprofile_start)*) dlsym(
handle, "hsa_ven_amd_aqlprofile_start");
api->hsa_ven_amd_aqlprofile_stop = (decltype(::hsa_ven_amd_aqlprofile_stop)*) dlsym(
handle, "hsa_ven_amd_aqlprofile_stop");
#ifdef AQLPROF_NEW_API
api->hsa_ven_amd_aqlprofile_read = (decltype(::hsa_ven_amd_aqlprofile_read)*) dlsym(
handle, "hsa_ven_amd_aqlprofile_read");
#endif
api->hsa_ven_amd_aqlprofile_legacy_get_pm4 =
(decltype(::hsa_ven_amd_aqlprofile_legacy_get_pm4)*) dlsym(
handle, "hsa_ven_amd_aqlprofile_legacy_get_pm4");
api->hsa_ven_amd_aqlprofile_get_info =
(decltype(::hsa_ven_amd_aqlprofile_get_info)*) dlsym(
handle, "hsa_ven_amd_aqlprofile_get_info");
api->hsa_ven_amd_aqlprofile_iterate_data =
(decltype(::hsa_ven_amd_aqlprofile_iterate_data)*) dlsym(
handle, "hsa_ven_amd_aqlprofile_iterate_data");
return HSA_STATUS_SUCCESS;
}
// Add system agent info
const AgentInfo*
HsaRsrcFactory::AddAgentInfo(const hsa_agent_t agent)
{
// Determine if device is a Gpu agent
hsa_status_t status;
AgentInfo* agent_info = nullptr;
hsa_device_type_t type;
status = hsa_api_.hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type);
CHECK_STATUS("Error Calling hsa_agent_get_info", status);
if(type == HSA_DEVICE_TYPE_CPU)
{
agent_info = new AgentInfo{};
agent_info->dev_id = agent;
agent_info->dev_type = HSA_DEVICE_TYPE_CPU;
agent_info->dev_index = cpu_list_.size();
status = hsa_api_.hsa_amd_agent_iterate_memory_pools(agent, FindStandardPool,
&agent_info->cpu_pool);
if((status == HSA_STATUS_INFO_BREAK) && (cpu_pool_ == nullptr))
cpu_pool_ = &agent_info->cpu_pool;
status = hsa_api_.hsa_amd_agent_iterate_memory_pools(agent, FindKernArgPool,
&agent_info->kern_arg_pool);
if((status == HSA_STATUS_INFO_BREAK) && (kern_arg_pool_ == nullptr))
kern_arg_pool_ = &agent_info->kern_arg_pool;
agent_info->gpu_pool = {};
cpu_list_.push_back(agent_info);
cpu_agents_.push_back(agent);
}
if(type == HSA_DEVICE_TYPE_GPU)
{
agent_info = new AgentInfo{};
agent_info->dev_id = agent;
agent_info->dev_type = HSA_DEVICE_TYPE_GPU;
hsa_api_.hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, agent_info->name);
strncpy(agent_info->gfxip, agent_info->name, 4);
agent_info->gfxip[4] = '\0';
hsa_api_.hsa_agent_get_info(agent, HSA_AGENT_INFO_WAVEFRONT_SIZE,
&agent_info->max_wave_size);
hsa_api_.hsa_agent_get_info(agent, HSA_AGENT_INFO_QUEUE_MAX_SIZE,
&agent_info->max_queue_size);
hsa_api_.hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_info->profile);
agent_info->is_apu = (agent_info->profile == HSA_PROFILE_FULL) ? true : false;
hsa_api_.hsa_agent_get_info(
agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT),
&agent_info->cu_num);
hsa_api_.hsa_agent_get_info(
agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU),
&agent_info->waves_per_cu);
hsa_api_.hsa_agent_get_info(
agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SIMDS_PER_CU),
&agent_info->simds_per_cu);
hsa_api_.hsa_agent_get_info(
agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SHADER_ENGINES),
&agent_info->se_num);
hsa_api_.hsa_agent_get_info(
agent,
static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SHADER_ARRAYS_PER_SE),
&agent_info->shader_arrays_per_se);
agent_info->cpu_pool = {};
agent_info->kern_arg_pool = {};
status = hsa_api_.hsa_amd_agent_iterate_memory_pools(agent, FindStandardPool,
&agent_info->gpu_pool);
CHECK_ITER_STATUS("hsa_amd_agent_iterate_memory_pools(gpu pool)", status);
// GFX8 and GFX9 SGPR/VGPR block sizes
agent_info->sgpr_block_dflt = (strcmp(agent_info->gfxip, "gfx8") == 0) ? 1 : 2;
agent_info->sgpr_block_size = 8;
agent_info->vgpr_block_size = 4;
// Set GPU index
agent_info->dev_index = gpu_list_.size();
gpu_list_.push_back(agent_info);
gpu_agents_.push_back(agent);
}
if(agent_info) agent_map_[agent.handle] = agent_info;
return agent_info;
}
// Return systen agent info
const AgentInfo*
HsaRsrcFactory::GetAgentInfo(const hsa_agent_t agent)
{
const AgentInfo* agent_info = nullptr;
auto it = agent_map_.find(agent.handle);
if(it != agent_map_.end())
{
agent_info = it->second;
}
return agent_info;
}
// Get the count of Hsa Gpu Agents available on the platform
//
// @return uint32_t Number of Gpu agents on platform
//
uint32_t
HsaRsrcFactory::GetCountOfGpuAgents()
{
return uint32_t(gpu_list_.size());
}
// Get the count of Hsa Cpu Agents available on the platform
//
// @return uint32_t Number of Cpu agents on platform
//
uint32_t
HsaRsrcFactory::GetCountOfCpuAgents()
{
return uint32_t(cpu_list_.size());
}
// Get the AgentInfo handle of a Gpu device
//
// @param idx Gpu Agent at specified index
//
// @param agent_info Output parameter updated with AgentInfo
//
// @return bool true if successful, false otherwise
//
bool
HsaRsrcFactory::GetGpuAgentInfo(uint32_t idx, const AgentInfo** agent_info)
{
// Determine if request is valid
uint32_t size = uint32_t(gpu_list_.size());
if(idx >= size)
{
return false;
}
// Copy AgentInfo from specified index
*agent_info = gpu_list_[idx];
return true;
}
// Get the AgentInfo handle of a Cpu device
//
// @param idx Cpu Agent at specified index
//
// @param agent_info Output parameter updated with AgentInfo
//
// @return bool true if successful, false otherwise
//
bool
HsaRsrcFactory::GetCpuAgentInfo(uint32_t idx, const AgentInfo** agent_info)
{
// Determine if request is valid
uint32_t size = uint32_t(cpu_list_.size());
if(idx >= size)
{
return false;
}
// Copy AgentInfo from specified index
*agent_info = cpu_list_[idx];
return true;
}
// Create a Queue object and return its handle. The queue object is expected
// to support user requested number of Aql dispatch packets.
//
// @param agent_info Gpu Agent on which to create a queue object
//
// @param num_Pkts Number of packets to be held by queue
//
// @param queue Output parameter updated with handle of queue object
//
// @return bool true if successful, false otherwise
//
bool
HsaRsrcFactory::CreateQueue(const AgentInfo* agent_info, uint32_t num_pkts,
hsa_queue_t** queue)
{
hsa_status_t status;
status = hsa_api_.hsa_queue_create(agent_info->dev_id, num_pkts, HSA_QUEUE_TYPE_MULTI,
nullptr, nullptr, UINT32_MAX, UINT32_MAX, queue);
return (status == HSA_STATUS_SUCCESS);
}
// Create a Signal object and return its handle.
// @param value Initial value of signal object
// @param signal Output parameter updated with handle of signal object
// @return bool true if successful, false otherwise
bool
HsaRsrcFactory::CreateSignal(uint32_t value, hsa_signal_t* signal)
{
hsa_status_t status;
status = hsa_api_.hsa_signal_create(value, 0, nullptr, signal);
return (status == HSA_STATUS_SUCCESS);
}
// Allocate memory for use by a kernel of specified size in specified
// agent's memory region.
// @param agent_info Agent from whose memory region to allocate
// @param size Size of memory in terms of bytes
// @return uint8_t* Pointer to buffer, null if allocation fails.
uint8_t*
HsaRsrcFactory::AllocateLocalMemory(const AgentInfo* agent_info, size_t size)
{
hsa_status_t status = HSA_STATUS_ERROR;
uint8_t* buffer = nullptr;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = hsa_api_.hsa_amd_memory_pool_allocate(agent_info->gpu_pool, size, 0,
reinterpret_cast<void**>(&buffer));
uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : nullptr;
return ptr;
}
// Allocate memory to pass kernel parameters.
// Memory is alocated accessible for all CPU agents and for GPU given by AgentInfo
// parameter.
// @param agent_info Agent from whose memory region to allocate
// @param size Size of memory in terms of bytes
// @return uint8_t* Pointer to buffer, null if allocation fails.
uint8_t*
HsaRsrcFactory::AllocateKernArgMemory(const AgentInfo* agent_info, size_t size)
{
hsa_status_t status = HSA_STATUS_ERROR;
uint8_t* buffer = nullptr;
if(!cpu_agents_.empty())
{
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
status = hsa_api_.hsa_amd_memory_pool_allocate(*kern_arg_pool_, size, 0,
reinterpret_cast<void**>(&buffer));
// Both the CPU and GPU can access the kernel arguments
if(status == HSA_STATUS_SUCCESS)
{
hsa_agent_t ag_list[1] = { agent_info->dev_id };
status = hsa_api_.hsa_amd_agents_allow_access(1, ag_list, nullptr, buffer);
}
}
uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : nullptr;
return ptr;
}
// Allocate system memory accessible by both CPU and GPU
// @param agent_info Agent from whose memory region to allocate
// @param size Size of memory in terms of bytes
// @return uint8_t* Pointer to buffer, null if allocation fails.
uint8_t*
HsaRsrcFactory::AllocateSysMemory(const AgentInfo* agent_info, size_t size)
{
hsa_status_t status = HSA_STATUS_ERROR;
uint8_t* buffer = nullptr;
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
if(!cpu_agents_.empty())
{
status = hsa_api_.hsa_amd_memory_pool_allocate(*cpu_pool_, size, 0,
reinterpret_cast<void**>(&buffer));
// Both the CPU and GPU can access the memory
if(status == HSA_STATUS_SUCCESS)
{
hsa_agent_t ag_list[1] = { agent_info->dev_id };
status = hsa_api_.hsa_amd_agents_allow_access(1, ag_list, nullptr, buffer);
}
}
uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : nullptr;
return ptr;
}
// Allocate memory for command buffer.
// @param agent_info Agent from whose memory region to allocate
// @param size Size of memory in terms of bytes
// @return uint8_t* Pointer to buffer, null if allocation fails.
uint8_t*
HsaRsrcFactory::AllocateCmdMemory(const AgentInfo* agent_info, size_t size)
{
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
uint8_t* ptr = (agent_info->is_apu && CMD_MEMORY_MMAP)
? reinterpret_cast<uint8_t*>(
mmap(nullptr, size, PROT_READ | PROT_WRITE | PROT_EXEC,
MAP_SHARED | MAP_ANONYMOUS, 0, 0))
: AllocateSysMemory(agent_info, size);
return ptr;
}
// Wait signal
void
HsaRsrcFactory::SignalWait(const hsa_signal_t& signal) const
{
while(1)
{
const hsa_signal_value_t signal_value = hsa_api_.hsa_signal_wait_scacquire(
signal, HSA_SIGNAL_CONDITION_LT, 1, timeout_, HSA_WAIT_STATE_BLOCKED);
if(signal_value == 0)
{
break;
}
else
{
CHECK_STATUS("hsa_signal_wait_scacquire()", HSA_STATUS_ERROR);
}
}
}
// Wait signal with signal value restore
void
HsaRsrcFactory::SignalWaitRestore(const hsa_signal_t& signal,
const hsa_signal_value_t& signal_value) const
{
SignalWait(signal);
hsa_api_.hsa_signal_store_relaxed(const_cast<hsa_signal_t&>(signal), signal_value);
}
// Copy data from GPU to host memory
bool
HsaRsrcFactory::Memcpy(const hsa_agent_t& agent, void* dst, const void* src, size_t size)
{
hsa_status_t status = HSA_STATUS_ERROR;
if(!cpu_agents_.empty())
{
hsa_signal_t s = {};
status = hsa_api_.hsa_signal_create(1, 0, nullptr, &s);
CHECK_STATUS("hsa_signal_create()", status);
status = hsa_api_.hsa_amd_memory_async_copy(dst, cpu_agents_[0], src, agent, size,
0, nullptr, s);
CHECK_STATUS("hsa_amd_memory_async_copy()", status);
SignalWait(s);
status = hsa_api_.hsa_signal_destroy(s);
CHECK_STATUS("hsa_signal_destroy()", status);
}
return (status == HSA_STATUS_SUCCESS);
}
bool
HsaRsrcFactory::Memcpy(const AgentInfo* agent_info, void* dst, const void* src,
size_t size)
{
return Memcpy(agent_info->dev_id, dst, src, size);
}
// Memory free method
bool
HsaRsrcFactory::FreeMemory(void* ptr)
{
const hsa_status_t status = hsa_memory_free(ptr);
CHECK_STATUS("hsa_memory_free", status);
return (status == HSA_STATUS_SUCCESS);
}
// Loads an Assembled Brig file and Finalizes it into Device Isa
// @param agent_info Gpu device for which to finalize
// @param brig_path File path of the Assembled Brig file
// @param kernel_name Name of the kernel to finalize
// @param code_desc Handle of finalized Code Descriptor that could
// be used to submit for execution
// @return bool true if successful, false otherwise
bool
HsaRsrcFactory::LoadAndFinalize(const AgentInfo* agent_info, const char* brig_path,
const char* kernel_name, hsa_executable_t* executable,
hsa_executable_symbol_t* code_desc)
{
hsa_status_t status = HSA_STATUS_ERROR;
// Build the code object filename
std::string filename(brig_path);
std::clog << "Code object filename: " << filename << std::endl;
// Open the file containing code object
hsa_file_t file_handle = open(filename.c_str(), O_RDONLY);
if(file_handle == -1)
{
std::cerr << "Error: failed to load '" << filename << "'" << std::endl;
assert(false);
return false;
}
// Create code object reader
hsa_code_object_reader_t code_obj_rdr = { 0 };
status = hsa_api_.hsa_code_object_reader_create_from_file(file_handle, &code_obj_rdr);
if(status != HSA_STATUS_SUCCESS)
{
std::cerr << "Failed to create code object reader '" << filename << "'"
<< std::endl;
return false;
}
// Create executable.
status = hsa_api_.hsa_executable_create_alt(
HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr, executable);
CHECK_STATUS("Error in creating executable object", status);
// Load code object.
status = hsa_api_.hsa_executable_load_agent_code_object(
*executable, agent_info->dev_id, code_obj_rdr, nullptr, nullptr);
CHECK_STATUS("Error in loading executable object", status);
// Freeze executable.
status = hsa_api_.hsa_executable_freeze(*executable, "");
CHECK_STATUS("Error in freezing executable object", status);
// Get symbol handle.
hsa_executable_symbol_t kernelSymbol;
status = hsa_api_.hsa_executable_get_symbol(*executable, nullptr, kernel_name,
agent_info->dev_id, 0, &kernelSymbol);
CHECK_STATUS("Error in looking up kernel symbol", status);
// Update output parameter
*code_desc = kernelSymbol;
return true;
}
// Print the various fields of Hsa Gpu Agents
bool
HsaRsrcFactory::PrintGpuAgents(const std::string&)
{
std::cout << std::flush;
// std::clog << header << " :" << std::endl;
char key[1024], value[1024];
const AgentInfo* agent_info;
int size = uint32_t(gpu_list_.size());
for(int idx = 0; idx < size; idx++)
{
agent_info = gpu_list_[idx];
/* std::clog << "> agent[" << idx << "] :" << std::endl;
std::clog << ">> Name : " << agent_info->name << std::endl;
std::clog << ">> APU : " << agent_info->is_apu << std::endl;
std::clog << ">> HSAIL profile : " << agent_info->profile << std::endl;
std::clog << ">> Max Wave Size : " << agent_info->max_wave_size << std::endl;
std::clog << ">> Max Queue Size : " << agent_info->max_queue_size <<
std::endl; std::clog << ">> CU number : " << agent_info->cu_num << std::endl;
std::clog << ">> Waves per CU : " << agent_info->waves_per_cu << std::endl;
std::clog << ">> SIMDs per CU : " << agent_info->simds_per_cu << std::endl;
std::clog << ">> SE number : " << agent_info->se_num << std::endl;
std::clog << ">> Shader Arrays per SE : " << agent_info->shader_arrays_per_se
<< std::endl;
*/
#define OMNITRACE_METADATA(...) tim::manager::add_metadata(__VA_ARGS__)
sprintf(key, "ROCM_AGENT_%d_NAME", idx);
sprintf(value, "%s", agent_info->name);
OMNITRACE_METADATA(key, value);
sprintf(key, "ROCM_AGENT_%d_IS_APU", idx);
sprintf(value, "%d", static_cast<int>(agent_info->is_apu));
OMNITRACE_METADATA(key, value);
sprintf(key, "ROCM_AGENT_%d_HSA_PROFILE", idx);
sprintf(value, "%d", agent_info->profile);
OMNITRACE_METADATA(key, value);
sprintf(key, "ROCM_AGENT_%d_MAX_WAVE_SIZE", idx);
sprintf(value, "%d", agent_info->max_wave_size);
OMNITRACE_METADATA(key, value);
sprintf(key, "ROCM_AGENT_%d_MAX_QUEUE_SIZE", idx);
sprintf(value, "%d", agent_info->max_queue_size);
OMNITRACE_METADATA(key, value);
sprintf(key, "ROCM_AGENT_%d_CU_NUMBER", idx);
sprintf(value, "%d", agent_info->cu_num);
OMNITRACE_METADATA(key, value);
sprintf(key, "ROCM_AGENT_%d_WAVES_PER_CU", idx);
sprintf(value, "%d", agent_info->waves_per_cu);
OMNITRACE_METADATA(key, value);
sprintf(key, "ROCM_AGENT_%d_SIMDs_PER_CU", idx);
sprintf(value, "%d", agent_info->simds_per_cu);
OMNITRACE_METADATA(key, value);
sprintf(key, "ROCM_AGENT_%d_SE_NUMBER", idx);
sprintf(value, "%d", agent_info->se_num);
OMNITRACE_METADATA(key, value);
sprintf(key, "ROCM_AGENT_%d_SHADER_ARRAYS_PER_SE", idx);
sprintf(value, "%d", agent_info->shader_arrays_per_se);
OMNITRACE_METADATA(key, value);
}
return true;
}
uint64_t
HsaRsrcFactory::Submit(hsa_queue_t* queue, const void* packet)
{
const uint32_t slot_size_b = CMD_SLOT_SIZE_B;
// adevance command queue
const uint64_t write_idx = hsa_api_.hsa_queue_load_write_index_relaxed(queue);
hsa_api_.hsa_queue_store_write_index_relaxed(queue, write_idx + 1);
while((write_idx - hsa_api_.hsa_queue_load_read_index_relaxed(queue)) >= queue->size)
{
sched_yield();
}
uint32_t slot_idx = (uint32_t)(write_idx % queue->size);
uint32_t* queue_slot = reinterpret_cast<uint32_t*>((uintptr_t)(queue->base_address) +
(slot_idx * slot_size_b));
const uint32_t* slot_data = reinterpret_cast<const uint32_t*>(packet);
// Copy buffered commands into the queue slot.
// Overwrite the AQL invalid header (first dword) last.
// This prevents the slot from being read until it's fully written.
memcpy(&queue_slot[1], &slot_data[1], slot_size_b - sizeof(uint32_t));
std::atomic<uint32_t>* header_atomic_ptr =
reinterpret_cast<std::atomic<uint32_t>*>(&queue_slot[0]);
header_atomic_ptr->store(slot_data[0], std::memory_order_release);
// ringdoor bell
hsa_api_.hsa_signal_store_relaxed(queue->doorbell_signal, write_idx);
return write_idx;
}
uint64_t
HsaRsrcFactory::Submit(hsa_queue_t* queue, const void* packet, size_t size_bytes)
{
const uint32_t slot_size_b = CMD_SLOT_SIZE_B;
if((size_bytes & (slot_size_b - 1)) != 0)
{
fprintf(stderr, "HsaRsrcFactory::Submit: Bad packet size %zx\n", size_bytes);
abort();
}
const char* begin = reinterpret_cast<const char*>(packet);
const char* end = begin + size_bytes;
uint64_t write_idx = 0;
for(const char* ptr = begin; ptr < end; ptr += slot_size_b)
{
write_idx = Submit(queue, ptr);
}
return write_idx;
}
const char*
HsaRsrcFactory::GetKernelName(uint64_t addr)
{
std::lock_guard<mutex_t> lck(mutex_);
const auto it = symbols_map_->find(addr);
if(it == symbols_map_->end())
{
fprintf(stderr, "HsaRsrcFactory::kernel addr (0x%lx) is not found\n", addr);
abort();
}
return strdup(it->second);
}
void
HsaRsrcFactory::EnableExecutableTracking(HsaApiTable* table)
{
std::lock_guard<mutex_t> lck(mutex_);
executable_tracking_on_ = true;
table->core_->hsa_executable_freeze_fn = hsa_executable_freeze_interceptor;
}
hsa_status_t
HsaRsrcFactory::executable_symbols_cb(hsa_executable_t /*exec*/,
hsa_executable_symbol_t symbol, void* /*data*/)
{
hsa_symbol_kind_t value = (hsa_symbol_kind_t) 0;
hsa_status_t status = hsa_api_.hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &value);
CHECK_STATUS("Error in getting symbol info", status);
if(value == HSA_SYMBOL_KIND_KERNEL)
{
uint64_t addr = 0;
uint32_t len = 0;
status = hsa_api_.hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &addr);
CHECK_STATUS("Error in getting kernel object", status);
status = hsa_api_.hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &len);
CHECK_STATUS("Error in getting name len", status);
char* name = new char[len + 1];
status = hsa_api_.hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
CHECK_STATUS("Error in getting kernel name", status);
name[len] = 0;
auto ret = symbols_map_->insert({ addr, name });
if(ret.second == false)
{
delete[] ret.first->second;
ret.first->second = name;
}
}
return HSA_STATUS_SUCCESS;
}
hsa_status_t
HsaRsrcFactory::hsa_executable_freeze_interceptor(hsa_executable_t executable,
const char* options)
{
std::lock_guard<mutex_t> lck(mutex_);
if(symbols_map_ == nullptr) symbols_map_ = new symbols_map_t;
hsa_status_t status = hsa_api_.hsa_executable_iterate_symbols(
executable, executable_symbols_cb, nullptr);
CHECK_STATUS("Error in iterating executable symbols", status);
return hsa_api_.hsa_executable_freeze(executable, options);
;
}
std::atomic<HsaRsrcFactory*> HsaRsrcFactory::instance_{};
HsaRsrcFactory::mutex_t HsaRsrcFactory::mutex_;
HsaRsrcFactory::timestamp_t HsaRsrcFactory::timeout_ns_ = HsaTimer::TIMESTAMP_MAX;
hsa_pfn_t HsaRsrcFactory::hsa_api_{};
bool HsaRsrcFactory::executable_tracking_on_ = false;
HsaRsrcFactory::symbols_map_t* HsaRsrcFactory::symbols_map_ = nullptr;
projects/rocprofiler-systems/source/lib/omnitrace/library/rocprofiler/hsa_rsrc_factory.hpp
+470
Parādīt failu
@@ -0,0 +1,470 @@
// MIT License
//
// Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in all
// copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
// SOFTWARE.
#pragma once
#define AMD_INTERNAL_BUILD 1
#include <hsa.h>
#include <hsa_api_trace.h>
#include <hsa_ext_amd.h>
#include <hsa_ext_finalize.h>
#include <hsa_ven_amd_aqlprofile.h>
#include <hsa_ven_amd_loader.h>
#include <atomic>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <iostream>
#include <map>
#include <mutex>
#include <string>
#include <vector>
#define HSA_ARGUMENT_ALIGN_BYTES 16
#define HSA_QUEUE_ALIGN_BYTES 64
#define HSA_PACKET_ALIGN_BYTES 64
#define HSA_MESSAGE_LENGTH 4096
#define CHECK_STATUS(msg, status) \
do \
{ \
if((status) != HSA_STATUS_SUCCESS) \
{ \
const char* emsg = 0; \
hsa_status_string(status, &emsg); \
char _buffer[HSA_MESSAGE_LENGTH]; \
snprintf(_buffer, HSA_MESSAGE_LENGTH - 1, "%s: %s", msg, \
emsg ? emsg : "<unknown error>"); \
throw std::runtime_error(_buffer); \
} \
} while(0)
#define CHECK_ITER_STATUS(msg, status) \
do \
{ \
if((status) != HSA_STATUS_INFO_BREAK) \
{ \
const char* emsg = 0; \
hsa_status_string(status, &emsg); \
char _buffer[HSA_MESSAGE_LENGTH]; \
snprintf(_buffer, HSA_MESSAGE_LENGTH - 1, "%s: %s", msg, \
emsg ? emsg : "<unknown error>"); \
throw std::runtime_error(_buffer); \
} \
} while(0)
static const size_t MEM_PAGE_BYTES = 0x1000;
static const size_t MEM_PAGE_MASK = MEM_PAGE_BYTES - 1;
typedef decltype(hsa_agent_t::handle) hsa_agent_handle_t;
struct hsa_pfn_t
{
decltype(::hsa_init)* hsa_init;
decltype(::hsa_shut_down)* hsa_shut_down;
decltype(::hsa_agent_get_info)* hsa_agent_get_info;
decltype(::hsa_iterate_agents)* hsa_iterate_agents;
decltype(::hsa_queue_create)* hsa_queue_create;
decltype(::hsa_queue_destroy)* hsa_queue_destroy;
decltype(::hsa_queue_load_write_index_relaxed)* hsa_queue_load_write_index_relaxed;
decltype(::hsa_queue_store_write_index_relaxed)* hsa_queue_store_write_index_relaxed;
decltype(::hsa_queue_load_read_index_relaxed)* hsa_queue_load_read_index_relaxed;
decltype(::hsa_signal_create)* hsa_signal_create;
decltype(::hsa_signal_destroy)* hsa_signal_destroy;
decltype(::hsa_signal_load_relaxed)* hsa_signal_load_relaxed;
decltype(::hsa_signal_store_relaxed)* hsa_signal_store_relaxed;
decltype(::hsa_signal_wait_scacquire)* hsa_signal_wait_scacquire;
decltype(::hsa_signal_store_screlease)* hsa_signal_store_screlease;
decltype(::hsa_code_object_reader_create_from_file)*
hsa_code_object_reader_create_from_file;
decltype(::hsa_executable_create_alt)* hsa_executable_create_alt;
decltype(
::hsa_executable_load_agent_code_object)* hsa_executable_load_agent_code_object;
decltype(::hsa_executable_freeze)* hsa_executable_freeze;
decltype(::hsa_executable_get_symbol)* hsa_executable_get_symbol;
decltype(::hsa_executable_symbol_get_info)* hsa_executable_symbol_get_info;
decltype(::hsa_executable_iterate_symbols)* hsa_executable_iterate_symbols;
decltype(::hsa_system_get_info)* hsa_system_get_info;
decltype(
::hsa_system_get_major_extension_table)* hsa_system_get_major_extension_table;
decltype(::hsa_amd_agent_iterate_memory_pools)* hsa_amd_agent_iterate_memory_pools;
decltype(::hsa_amd_memory_pool_get_info)* hsa_amd_memory_pool_get_info;
decltype(::hsa_amd_memory_pool_allocate)* hsa_amd_memory_pool_allocate;
decltype(::hsa_amd_agents_allow_access)* hsa_amd_agents_allow_access;
decltype(::hsa_amd_memory_async_copy)* hsa_amd_memory_async_copy;
decltype(::hsa_amd_signal_async_handler)* hsa_amd_signal_async_handler;
decltype(
::hsa_amd_profiling_set_profiler_enabled)* hsa_amd_profiling_set_profiler_enabled;
decltype(
::hsa_amd_profiling_get_async_copy_time)* hsa_amd_profiling_get_async_copy_time;
decltype(::hsa_amd_profiling_get_dispatch_time)* hsa_amd_profiling_get_dispatch_time;
};
// Encapsulates information about a Hsa Agent such as its
// handle, name, max queue size, max wavefront size, etc.
struct AgentInfo
{
// Handle of Agent
hsa_agent_t dev_id;
// Agent type - Cpu = 0, Gpu = 1 or Dsp = 2
uint32_t dev_type;
// APU flag
bool is_apu;
// Agent system index
uint32_t dev_index;
// GFXIP name
char gfxip[64];
// Name of Agent whose length is less than 64
char name[64];
// Max size of Wavefront size
uint32_t max_wave_size;
// Max size of Queue buffer
uint32_t max_queue_size;
// Hsail profile supported by agent
hsa_profile_t profile;
// CPU/GPU/kern-arg memory pools
hsa_amd_memory_pool_t cpu_pool;
hsa_amd_memory_pool_t gpu_pool;
hsa_amd_memory_pool_t kern_arg_pool;
// The number of compute unit available in the agent.
uint32_t cu_num;
// Maximum number of waves possible in a Compute Unit.
uint32_t waves_per_cu;
// Number of SIMD's per compute unit CU
uint32_t simds_per_cu;
// Number of Shader Engines (SE) in Gpu
uint32_t se_num;
// Number of Shader Arrays Per Shader Engines in Gpu
uint32_t shader_arrays_per_se;
// SGPR/VGPR block sizes
uint32_t sgpr_block_dflt;
uint32_t sgpr_block_size;
uint32_t vgpr_block_size;
};
// HSA timer class
// Provides current HSA timestampa and system-clock/ns conversion API
class HsaTimer
{
public:
typedef uint64_t timestamp_t;
static const timestamp_t TIMESTAMP_MAX = UINT64_MAX;
typedef long double freq_t;
HsaTimer(const hsa_pfn_t* hsa_api)
: hsa_api_(hsa_api)
{
timestamp_t sysclock_hz = 0;
hsa_status_t status = hsa_api_->hsa_system_get_info(
HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &sysclock_hz);
CHECK_STATUS("hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY)", status);
sysclock_factor_ = (freq_t) 1000000000 / (freq_t) sysclock_hz;
}
// Methods for system-clock/ns conversion
timestamp_t sysclock_to_ns(const timestamp_t& sysclock) const
{
return timestamp_t((freq_t) sysclock * sysclock_factor_);
}
timestamp_t ns_to_sysclock(const timestamp_t& time) const
{
return timestamp_t((freq_t) time / sysclock_factor_);
}
// Return timestamp in 'ns'
timestamp_t timestamp_ns() const
{
timestamp_t sysclock;
hsa_status_t status =
hsa_api_->hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP, &sysclock);
CHECK_STATUS("hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP)", status);
return sysclock_to_ns(sysclock);
}
private:
// Timestamp frequency factor
freq_t sysclock_factor_;
// HSA API table
const hsa_pfn_t* const hsa_api_;
};
class HsaRsrcFactory
{
public:
static const size_t CMD_SLOT_SIZE_B = 0x40;
typedef std::recursive_mutex mutex_t;
typedef HsaTimer::timestamp_t timestamp_t;
static HsaRsrcFactory* Create(bool initialize_hsa = true)
{
std::lock_guard<mutex_t> lck(mutex_);
HsaRsrcFactory* obj = instance_.load(std::memory_order_relaxed);
if(obj == nullptr)
{
obj = new HsaRsrcFactory(initialize_hsa);
instance_.store(obj, std::memory_order_release);
}
return obj;
}
static HsaRsrcFactory& Instance()
{
HsaRsrcFactory* obj = instance_.load(std::memory_order_acquire);
if(obj == nullptr) obj = Create(false);
hsa_status_t status = (obj != nullptr) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR;
CHECK_STATUS("HsaRsrcFactory::Instance() failed", status);
return *obj;
}
static void Destroy()
{
std::lock_guard<mutex_t> lck(mutex_);
if(instance_) delete instance_.load();
instance_ = nullptr;
}
// Return system agent info
const AgentInfo* GetAgentInfo(const hsa_agent_t agent);
// Get the count of Hsa Gpu Agents available on the platform
// @return uint32_t Number of Gpu agents on platform
uint32_t GetCountOfGpuAgents();
// Get the count of Hsa Cpu Agents available on the platform
// @return uint32_t Number of Cpu agents on platform
uint32_t GetCountOfCpuAgents();
// Get the AgentInfo handle of a Gpu device
// @param idx Gpu Agent at specified index
// @param agent_info Output parameter updated with AgentInfo
// @return bool true if successful, false otherwise
bool GetGpuAgentInfo(uint32_t idx, const AgentInfo** agent_info);
// Get the AgentInfo handle of a Cpu device
// @param idx Cpu Agent at specified index
// @param agent_info Output parameter updated with AgentInfo
// @return bool true if successful, false otherwise
bool GetCpuAgentInfo(uint32_t idx, const AgentInfo** agent_info);
// Create a Queue object and return its handle. The queue object is expected
// to support user requested number of Aql dispatch packets.
// @param agent_info Gpu Agent on which to create a queue object
// @param num_Pkts Number of packets to be held by queue
// @param queue Output parameter updated with handle of queue object
// @return bool true if successful, false otherwise
bool CreateQueue(const AgentInfo* agent_info, uint32_t num_pkts, hsa_queue_t** queue);
// Create a Signal object and return its handle.
// @param value Initial value of signal object
// @param signal Output parameter updated with handle of signal object
// @return bool true if successful, false otherwise
bool CreateSignal(uint32_t value, hsa_signal_t* signal);
// Allocate local GPU memory
// @param agent_info Agent from whose memory region to allocate
// @param size Size of memory in terms of bytes
// @return uint8_t* Pointer to buffer, null if allocation fails.
uint8_t* AllocateLocalMemory(const AgentInfo* agent_info, size_t size);
// Allocate memory tp pass kernel parameters
// Memory is alocated accessible for all CPU agents and for GPU given by AgentInfo
// parameter.
// @param agent_info Agent from whose memory region to allocate
// @param size Size of memory in terms of bytes
// @return uint8_t* Pointer to buffer, null if allocation fails.
uint8_t* AllocateKernArgMemory(const AgentInfo* agent_info, size_t size);
// Allocate system memory accessible from both CPU and GPU
// Memory is alocated accessible to all CPU agents and AgentInfo parameter is ignored.
// @param agent_info Agent from whose memory region to allocate
// @param size Size of memory in terms of bytes
// @return uint8_t* Pointer to buffer, null if allocation fails.
uint8_t* AllocateSysMemory(const AgentInfo* agent_info, size_t size);
// Allocate memory for command buffer.
// @param agent_info Agent from whose memory region to allocate
// @param size Size of memory in terms of bytes
// @return uint8_t* Pointer to buffer, null if allocation fails.
uint8_t* AllocateCmdMemory(const AgentInfo* agent_info, size_t size);
// Wait signal
void SignalWait(const hsa_signal_t& signal) const;
// Wait signal with signal value restore
void SignalWaitRestore(const hsa_signal_t& signal,
const hsa_signal_value_t& signal_value) const;
// Copy data from GPU to host memory
bool Memcpy(const hsa_agent_t& agent, void* dst, const void* src, size_t size);
bool Memcpy(const AgentInfo* agent_info, void* dst, const void* src, size_t size);
// Memory free method
static bool FreeMemory(void* ptr);
// Loads an Assembled Brig file and Finalizes it into Device Isa
// @param agent_info Gpu device for which to finalize
// @param brig_path File path of the Assembled Brig file
// @param kernel_name Name of the kernel to finalize
// @param code_desc Handle of finalized Code Descriptor that could
// be used to submit for execution
// @return true if successful, false otherwise
bool LoadAndFinalize(const AgentInfo* agent_info, const char* brig_path,
const char* kernel_name, hsa_executable_t* hsa_exec,
hsa_executable_symbol_t* code_desc);
// Print the various fields of Hsa Gpu Agents
bool PrintGpuAgents(const std::string& header);
// Submit AQL packet to given queue
static uint64_t Submit(hsa_queue_t* queue, const void* packet);
static uint64_t Submit(hsa_queue_t* queue, const void* packet, size_t size_bytes);
// Enable executables loading tracking
static bool IsExecutableTracking() { return executable_tracking_on_; }
static void EnableExecutableTracking(HsaApiTable* table);
static const char* GetKernelName(uint64_t addr);
// Initialize HSA API table
void static InitHsaApiTable(HsaApiTable* table);
static const hsa_pfn_t* HsaApi() { return &hsa_api_; }
// Return AqlProfile API table
typedef hsa_ven_amd_aqlprofile_pfn_t aqlprofile_pfn_t;
const aqlprofile_pfn_t* AqlProfileApi() const { return &aqlprofile_api_; }
// Return Loader API table
const hsa_ven_amd_loader_1_00_pfn_t* LoaderApi() const { return &loader_api_; }
// Methods for system-clock/ns conversion and timestamp in 'ns'
timestamp_t SysclockToNs(const timestamp_t& sysclock) const
{
return timer_->sysclock_to_ns(sysclock);
}
timestamp_t NsToSysclock(const timestamp_t& time) const
{
return timer_->ns_to_sysclock(time);
}
timestamp_t TimestampNs() const { return timer_->timestamp_ns(); }
timestamp_t GetSysTimeout() const { return timeout_; }
static timestamp_t GetTimeoutNs() { return timeout_ns_; }
static void SetTimeoutNs(const timestamp_t& time)
{
std::lock_guard<mutex_t> lck(mutex_);
timeout_ns_ = time;
if(instance_ != nullptr)
Instance().timeout_ = Instance().timer_->ns_to_sysclock(time);
}
private:
// System agents iterating callback
static hsa_status_t GetHsaAgentsCallback(hsa_agent_t agent, void* data);
// Callback function to find and bind kernarg region of an agent
static hsa_status_t FindMemRegionsCallback(hsa_region_t region, void* data);
// Load AQL profile HSA extension library directly
static hsa_status_t LoadAqlProfileLib(aqlprofile_pfn_t* api);
// Constructor of the class. Will initialize the Hsa Runtime and
// query the system topology to get the list of Cpu and Gpu devices
explicit HsaRsrcFactory(bool initialize_hsa);
// Destructor of the class
~HsaRsrcFactory();
// Add an instance of AgentInfo representing a Hsa Gpu agent
const AgentInfo* AddAgentInfo(const hsa_agent_t agent);
// To mmap command buffer memory
static const bool CMD_MEMORY_MMAP = false;
// HSA was initialized
const bool initialize_hsa_;
static std::atomic<HsaRsrcFactory*> instance_;
static mutex_t mutex_;
// Used to maintain a list of Hsa Gpu Agent Info
std::vector<const AgentInfo*> gpu_list_;
std::vector<hsa_agent_t> gpu_agents_;
// Used to maintain a list of Hsa Cpu Agent Info
std::vector<const AgentInfo*> cpu_list_;
std::vector<hsa_agent_t> cpu_agents_;
// System agents map
std::map<hsa_agent_handle_t, const AgentInfo*> agent_map_;
// Executables loading tracking
typedef std::map<uint64_t, const char*> symbols_map_t;
static symbols_map_t* symbols_map_;
static bool executable_tracking_on_;
static hsa_status_t hsa_executable_freeze_interceptor(hsa_executable_t executable,
const char* options);
static hsa_status_t executable_symbols_cb(hsa_executable_t exec,
hsa_executable_symbol_t symbol, void* data);
// HSA runtime API table
static hsa_pfn_t hsa_api_;
// AqlProfile API table
aqlprofile_pfn_t aqlprofile_api_;
// Loader API table
hsa_ven_amd_loader_1_00_pfn_t loader_api_;
// System timeout, ns
static timestamp_t timeout_ns_;
// System timeout, sysclock
timestamp_t timeout_;
// HSA timer
HsaTimer* timer_;
// CPU/kern-arg memory pools
hsa_amd_memory_pool_t* cpu_pool_;
hsa_amd_memory_pool_t* kern_arg_pool_;
};
projects/rocprofiler-systems/source/lib/omnitrace/library/components/roctracer_callbacks.cpp → projects/rocprofiler-systems/source/lib/omnitrace/library/roctracer.cpp
+72 -193
Parādīt failu
@@ -20,7 +20,7 @@
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
// SOFTWARE.
#include "library/components/roctracer_callbacks.hpp"
#include "library/roctracer.hpp"
#include "library.hpp"
#include "library/config.hpp"
#include "library/critical_trace.hpp"
@@ -77,65 +77,6 @@ hip_api_string(hip_api_id_t id, const hip_api_data_t* data)
tim::consume_parameters(id, data);
#endif
}
//
int64_t
get_clock_skew()
{
static auto _use = tim::get_env("OMNITRACE_USE_ROCTRACER_CLOCK_SKEW", true);
static auto _v = []() {
namespace cpu = tim::cpu;
// synchronize timestamps
// We'll take a CPU timestamp before and after taking a GPU timestmp, then
// take the average of those two, hoping that it's roughly at the same time
// as the GPU timestamp.
static auto _cpu_now = []() {
cpu::fence();
return comp::wall_clock::record();
};
static auto _gpu_now = []() {
cpu::fence();
uint64_t _v = 0;
ROCTRACER_CALL(roctracer_get_timestamp(&_v));
return _v;
};
do
{
// warm up cache and allow for any static initialization
(void) _cpu_now();
(void) _gpu_now();
} while(false);
auto _compute = [](volatile uint64_t& _cpu_ts, volatile uint64_t& _gpu_ts) {
_cpu_ts = 0;
_gpu_ts = 0;
_cpu_ts += _cpu_now() / 2;
_gpu_ts += _gpu_now() / 1;
_cpu_ts += _cpu_now() / 2;
return static_cast<int64_t>(_cpu_ts) - static_cast<int64_t>(_gpu_ts);
};
constexpr int64_t _n = 10;
int64_t _cpu_ave = 0;
int64_t _gpu_ave = 0;
int64_t _diff = 0;
for(int64_t i = 0; i < _n; ++i)
{
volatile uint64_t _cpu_ts = 0;
volatile uint64_t _gpu_ts = 0;
_diff += _compute(_cpu_ts, _gpu_ts);
_cpu_ave += _cpu_ts / _n;
_gpu_ave += _gpu_ts / _n;
}
OMNITRACE_BASIC_VERBOSE(2, "CPU timestamp: %li\n", _cpu_ave);
OMNITRACE_BASIC_VERBOSE(2, "HIP timestamp: %li\n", _gpu_ave);
OMNITRACE_BASIC_VERBOSE(1, "CPU/HIP timestamp skew: %li (used: %s)\n", _diff,
_use ? "yes" : "no");
_diff /= _n;
return _diff;
}();
return (_use) ? _v : 0;
}
int&
get_current_device()
@@ -218,6 +159,66 @@ get_hip_activity_mutex(int64_t _tid = threading::get_id())
}
} // namespace
//
int64_t
get_clock_skew()
{
static auto _use = tim::get_env("OMNITRACE_USE_ROCTRACER_CLOCK_SKEW", true);
static auto _v = []() {
namespace cpu = tim::cpu;
// synchronize timestamps
// We'll take a CPU timestamp before and after taking a GPU timestmp, then
// take the average of those two, hoping that it's roughly at the same time
// as the GPU timestamp.
static auto _cpu_now = []() {
cpu::fence();
return comp::wall_clock::record();
};
static auto _gpu_now = []() {
cpu::fence();
uint64_t _v = 0;
ROCTRACER_CALL(roctracer_get_timestamp(&_v));
return _v;
};
do
{
// warm up cache and allow for any static initialization
(void) _cpu_now();
(void) _gpu_now();
} while(false);
auto _compute = [](volatile uint64_t& _cpu_ts, volatile uint64_t& _gpu_ts) {
_cpu_ts = 0;
_gpu_ts = 0;
_cpu_ts += _cpu_now() / 2;
_gpu_ts += _gpu_now() / 1;
_cpu_ts += _cpu_now() / 2;
return static_cast<int64_t>(_cpu_ts) - static_cast<int64_t>(_gpu_ts);
};
constexpr int64_t _n = 10;
int64_t _cpu_ave = 0;
int64_t _gpu_ave = 0;
int64_t _diff = 0;
for(int64_t i = 0; i < _n; ++i)
{
volatile uint64_t _cpu_ts = 0;
volatile uint64_t _gpu_ts = 0;
_diff += _compute(_cpu_ts, _gpu_ts);
_cpu_ave += _cpu_ts / _n;
_gpu_ave += _gpu_ts / _n;
}
OMNITRACE_BASIC_VERBOSE(2, "CPU timestamp: %li\n", _cpu_ave);
OMNITRACE_BASIC_VERBOSE(2, "HIP timestamp: %li\n", _gpu_ave);
OMNITRACE_BASIC_VERBOSE(1, "CPU/HIP timestamp skew: %li (used: %s)\n", _diff,
_use ? "yes" : "no");
_diff /= _n;
return _diff;
}();
return (_use) ? _v : 0;
}
// HSA API callback function
void
hsa_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg)
@@ -227,6 +228,15 @@ hsa_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void*
OMNITRACE_SCOPED_THREAD_STATE(ThreadState::Internal);
static thread_local std::once_flag _once{};
std::call_once(_once, []() {
if(threading::get_id() != 0)
{
sampling::block_signals();
threading::set_thread_name("roctracer.hsa");
}
});
(void) arg;
const hsa_api_data_t* data = reinterpret_cast<const hsa_api_data_t*>(callback_data);
OMNITRACE_CONDITIONAL_PRINT_F(
@@ -343,7 +353,7 @@ hsa_activity_callback(uint32_t op, activity_record_t* record, void* arg)
static thread_local std::once_flag _once{};
std::call_once(_once, []() {
sampling::block_signals();
threading::set_thread_name("omni.roctracer");
threading::set_thread_name("roctracer.hsa");
});
auto&& _protect = comp::roctracer::protect_flush_activity();
@@ -699,7 +709,7 @@ hip_activity_callback(const char* begin, const char* end, void*)
static thread_local std::once_flag _once{};
std::call_once(_once, []() {
sampling::block_signals();
threading::set_thread_name("omni.roctracer");
threading::set_thread_name("roctracer.hip");
});
auto&& _protect = comp::roctracer::protect_flush_activity();
@@ -883,134 +893,3 @@ roctracer_shutdown_routines()
return _v;
}
} // namespace omnitrace
#include "library/components/rocm_smi.hpp"
using namespace omnitrace;
// HSA-runtime tool on-load method
extern "C"
{
bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
const char* const* failed_tool_names) OMNITRACE_VISIBILITY("default");
void OnUnload() OMNITRACE_VISIBILITY("default");
bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
const char* const* failed_tool_names)
{
if(!tim::get_env("OMNITRACE_INIT_TOOLING", true)) return true;
if(!tim::settings::enabled()) return true;
roctracer_is_init() = true;
pthread_gotcha::push_enable_sampling_on_child_threads(false);
OMNITRACE_CONDITIONAL_BASIC_PRINT_F(get_debug_env() || get_verbose_env() > 0,
"\n");
tim::consume_parameters(table, runtime_version, failed_tool_count,
failed_tool_names);
if(!config::settings_are_configured() && get_state() < State::Active)
omnitrace_init_tooling_hidden();
OMNITRACE_SCOPED_THREAD_STATE(ThreadState::Internal);
static auto _setup = [=]() {
try
{
OMNITRACE_CONDITIONAL_BASIC_PRINT_F(get_debug() || get_verbose() > 1,
"setting up HSA...\n");
// const char* output_prefix = getenv("ROCP_OUTPUT_DIR");
const char* output_prefix = nullptr;
bool trace_hsa_api = get_trace_hsa_api();
// Enable HSA API callbacks/activity
if(trace_hsa_api)
{
std::vector<std::string> hsa_api_vec =
tim::delimit(get_trace_hsa_api_types());
// initialize HSA tracing
roctracer_set_properties(ACTIVITY_DOMAIN_HSA_API, (void*) table);
OMNITRACE_CONDITIONAL_BASIC_PRINT(get_debug() || get_verbose() > 1,
" HSA-trace(");
if(!hsa_api_vec.empty())
{
for(const auto& itr : hsa_api_vec)
{
uint32_t cid = HSA_API_ID_NUMBER;
const char* api = itr.c_str();
ROCTRACER_CALL(roctracer_op_code(ACTIVITY_DOMAIN_HSA_API, api,
&cid, nullptr));
ROCTRACER_CALL(roctracer_enable_op_callback(
ACTIVITY_DOMAIN_HSA_API, cid, hsa_api_callback, nullptr));
OMNITRACE_CONDITIONAL_BASIC_PRINT(
get_debug() || get_verbose() > 1, " %s", api);
}
}
else
{
ROCTRACER_CALL(roctracer_enable_domain_callback(
ACTIVITY_DOMAIN_HSA_API, hsa_api_callback, nullptr));
}
OMNITRACE_CONDITIONAL_BASIC_PRINT(get_debug() || get_verbose() > 1,
"\n");
}
bool trace_hsa_activity = get_trace_hsa_activity();
// Enable HSA GPU activity
if(trace_hsa_activity)
{
// initialize HSA tracing
::roctracer::hsa_ops_properties_t ops_properties{
table,
reinterpret_cast<activity_async_callback_t>(
hsa_activity_callback),
nullptr, output_prefix
};
roctracer_set_properties(ACTIVITY_DOMAIN_HSA_OPS, &ops_properties);
OMNITRACE_CONDITIONAL_BASIC_PRINT(get_debug() || get_verbose() > 1,
" HSA-activity-trace()\n");
ROCTRACER_CALL(roctracer_enable_op_activity(ACTIVITY_DOMAIN_HSA_OPS,
HSA_OP_ID_COPY));
}
} catch(std::exception& _e)
{
OMNITRACE_BASIC_PRINT("Exception was thrown in HSA setup: %s\n",
_e.what());
}
};
static auto _shutdown = []() {
OMNITRACE_DEBUG_F("roctracer_disable_domain_callback\n");
ROCTRACER_CALL(roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HSA_API));
OMNITRACE_DEBUG_F("roctracer_disable_op_activity\n");
ROCTRACER_CALL(
roctracer_disable_op_activity(ACTIVITY_DOMAIN_HSA_OPS, HSA_OP_ID_COPY));
};
(void) get_clock_skew();
comp::roctracer::add_setup("hsa", _setup);
comp::roctracer::add_shutdown("hsa", _shutdown);
rocm_smi::set_state(State::Active);
comp::roctracer::setup();
pthread_gotcha::pop_enable_sampling_on_child_threads();
return true;
}
// HSA-runtime on-unload method
void OnUnload()
{
OMNITRACE_DEBUG_F("\n");
rocm_smi::set_state(State::Finalized);
comp::roctracer::shutdown();
omnitrace_finalize_hidden();
}
}
projects/rocprofiler-systems/source/lib/omnitrace/library/components/roctracer_callbacks.hpp → projects/rocprofiler-systems/source/lib/omnitrace/library/roctracer.hpp
+3
Parādīt failu
@@ -77,6 +77,9 @@ roctracer_is_init();
bool&
roctracer_is_setup();
int64_t
get_clock_skew();
roctracer_functions_t&
roctracer_setup_routines();
projects/rocprofiler-systems/source/lib/omnitrace/library/runtime.cpp
+4 -2
Parādīt failu
@@ -157,7 +157,8 @@ std::unique_ptr<main_bundle_t>&
get_main_bundle()
{
static auto _v =
std::make_unique<main_bundle_t>("omnitrace", quirk::config<quirk::auto_start>{});
std::make_unique<main_bundle_t>(JOIN('/', "omnitrace/process", process::get_id()),
quirk::config<quirk::auto_start>{});
return _v;
}
@@ -166,7 +167,8 @@ get_gotcha_bundle()
{
static auto _v =
(setup_gotchas(), std::make_unique<gotcha_bundle_t>(
"omnitrace", quirk::config<quirk::auto_start>{}));
JOIN('/', "omnitrace/process", process::get_id()),
quirk::config<quirk::auto_start>{}));
return _v;
}
projects/rocprofiler-systems/source/lib/omnitrace/library/timemory.hpp
+13 -11
Parādīt failu
@@ -44,15 +44,17 @@
namespace omnitrace
{
namespace audit = tim::audit; // NOLINT
namespace comp = tim::component; // NOLINT
namespace quirk = tim::quirk; // NOLINT
namespace threading = tim::threading; // NOLINT
namespace scope = tim::scope; // NOLINT
namespace dmp = tim::dmp; // NOLINT
namespace process = tim::process; // NOLINT
namespace units = tim::units; // NOLINT
namespace trait = tim::trait; // NOLINT
namespace api = tim::api; // NOLINT
namespace operation = tim::operation; // NOLINT
namespace audit = ::tim::audit; // NOLINT
namespace comp = ::tim::component; // NOLINT
namespace quirk = ::tim::quirk; // NOLINT
namespace threading = ::tim::threading; // NOLINT
namespace scope = ::tim::scope; // NOLINT
namespace dmp = ::tim::dmp; // NOLINT
namespace process = ::tim::process; // NOLINT
namespace units = ::tim::units; // NOLINT
namespace trait = ::tim::trait; // NOLINT
namespace api = ::tim::api; // NOLINT
namespace operation = ::tim::operation; // NOLINT
using settings = ::tim::settings; // NOLINT
} // namespace omnitrace
projects/rocprofiler-systems/source/lib/omnitrace/library/tracing.hpp
+1 -2
Parādīt failu
@@ -91,11 +91,10 @@ inline void
thread_init()
{
static thread_local auto _thread_setup = []() {
auto _exe = get_exe_name();
if(threading::get_id() > 0)
threading::set_thread_name(JOIN(" ", "Thread", threading::get_id()).c_str());
thread_data<omnitrace_thread_bundle_t>::construct(
JOIN("", _exe, "/thread-", threading::get_id()),
JOIN('/', "omnitrace", process::get_id(), "thread", threading::get_id()),
quirk::config<quirk::auto_start>{});
get_interval_data()->reserve(512);
// save the hash maps