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

Integrated perfetto + roctracer (#5)

Browse Source 
- hosttrace library automatically collects and merges timestamps for HIP API calls and kernels with the host-side instrumentation
  - mostly eliminates the need for using external rocprof
- added thread_instruction_count in perfetto output
- increased hosttrace min_loop_address_range to 512
- disabled instrumenting functions with dynamic callsites by default
- miscellaneous cmake updates

* roctracer support

- fully integrated perfetto + roctracer outputs
- thread_instruction_count in perfetto
- increased min_loop_address_range to 512
- disabled instrumenting functions with dynamic callsites by default
- updated timemory submodule

* hosttrace_launch_compiler

- support for using an alternative compiler as needed via launch compiler
- elfio added as submodule (not currently used)
- miscellaneous cmake updates

* README update + host/device categories + misc

- timemory fix for TIMEMORY_ROCTRACER_ENABLED
- transpose fix

* papi_tuple_t -> papi_tot_ins

- minor fix to Findroctracer.cmake
This commit is contained in:
Jonathan R. Madsen
2021-09-06 22:23:24 -05:00
committed by GitHub
parent b6e65d69f6
commit a061b7947f
23 changed files with 1119 additions and 268 deletions
Download Patch File Download Diff File Expand all files Collapse all files
.gitmodules
+3
View File
@@ -4,3 +4,6 @@
[submodule "external/perfetto"]
path = external/perfetto
url = https://android.googlesource.com/platform/external/perfetto
[submodule "external/elfio"]
path = external/elfio
url = https://github.com/jrmadsen/ELFIO.git
CMakeLists.txt
+8
View File
@@ -33,6 +33,7 @@ add_option(CMAKE_INSTALL_RPATH_USE_LINK_PATH "Enable rpath to linked libraries"
add_option(HOSTTRACE_USE_CLANG_TIDY "Enable clang-tidy" OFF)
add_option(HOSTTRACE_USE_MPI "Enable MPI support" OFF)
add_option(HOSTTRACE_CUSTOM_DATA_SOURCE "Enable custom data source" OFF)
add_option(HOSTTRACE_USE_ROCTRACER "Enable roctracer support" ON)
# ensure synced
set(TIMEMORY_USE_MPI ${HOSTTRACE_USE_MPI} CACHE BOOL "Enable MPI support" FORCE)
@@ -54,6 +55,12 @@ add_library(hosttrace-library SHARED
${CMAKE_CURRENT_LIST_DIR}/include/library.hpp
${perfetto_DIR}/sdk/perfetto.cc)
if(HOSTTRACE_USE_ROCTRACER)
target_sources(hosttrace-library PRIVATE
${CMAKE_CURRENT_LIST_DIR}/include/roctracer.hpp
${CMAKE_CURRENT_LIST_DIR}/src/roctracer.cpp)
endif()
target_include_directories(hosttrace-library PRIVATE
${CMAKE_CURRENT_LIST_DIR}/include)
@@ -66,6 +73,7 @@ target_compile_definitions(hosttrace-library PRIVATE
target_link_libraries(hosttrace-library PRIVATE
hosttrace::hosttrace-threading
hosttrace::hosttrace-compile-options
hosttrace::hosttrace-roctracer
$<BUILD_INTERFACE:timemory::timemory-headers>
$<BUILD_INTERFACE:timemory::timemory-gotcha>
$<BUILD_INTERFACE:timemory::timemory-cxx-shared>
README.md
+69 -19
View File
@@ -9,7 +9,7 @@
The easiest way to install Dyninst is via spack
```console
```shell
git clone https://github.com/spack/spack.git
source ./spack/share/spack/setup-env.sh
spack compiler find
@@ -22,19 +22,19 @@ spack load -r dyninst
Julia is available via Linux package managers or may be available via a module. Debian-based distributions such as Ubuntu can run (as a super-user):
```console
```shell
apt-get install julia
```
Once Julia is installed, install the necessary packages (this operation only needs to be performed once):
```console
```shell
julia -e 'using Pkg; for name in ["JSON", "DataFrames", "Dates", "CSV", "Chain", "PrettyTables"]; Pkg.add(name); end'
```
## Installing hosttrace
```console
```shell
HOSTTRACE_ROOT=${HOME}/sw/hosttrace
git clone https://github.com/AARInternal/hosttrace-dyninst.git
cmake -B build-hosttrace -DHOSTTRACE_USE_MPI=ON -DCMAKE_INSTALL_PREFIX=${HOSTTRACE_ROOT} hosttrace-dyninst
@@ -44,27 +44,43 @@ export PATH=${HOSTTRACE_ROOT}/bin:${PATH}
export LD_LIBRARY_PATH=${HOSTTRACE_ROOT}/lib64:${HOSTTRACE_ROOT}/lib:${LD_LIBRARY_PATH}
```
## Using hosttrace
## Using Hosttrace Executable
```console
```shell
hosttrace --help
hosttrace <hosttrace-options> -- <exe-or-library> <exe-options>
```
### Example Instrumentation
## Hosttrace Library Environment Settings
| Environment Variable | Default Value | Description |
|-----------------------------|-------------------------------|----------------------------------------------------------------------------------|
| `HOSTTRACE_DEBUG` | `false` | Enable debugging statements |
| `HOSTTRACE_USE_PERFETTO` | `true` | Collect profiling data via perfetto |
| `HOSTTRACE_USE_TIMEMORY` | `false` | Collection profiling data via timemory |
| `HOSTTRACE_SAMPLE_RATE` | `1` | Invoke perfetto and/or timemory once every N function calls |
| `HOSTTRACE_USE_MPI` | `true` | Label perfetto output files via rank instead of PID |
| `HOSTTRACE_OUTPUT_FILE` | `perfetto-trace.%rank%.proto` | Output file for perfetto (may use `%pid`) |
| `HOSTTRACE_BACKEND` | `"inprocess"` | Configure perfetto to use either "inprocess" data management, "system", or "all" |
| `HOSTTRACE_COMPONENTS` | `"wall_clock"` | Timemory components to activate when enabled |
| `HOSTTRACE_SHMEM_SIZE_HINT` | `40960` | Hint for perfetto shared memory buffer |
| `HOSTTRACE_BUFFER_SIZE_KB` | `1024000` | Maximum amount of memory perfetto will use to collect data in-process |
| `TIMEMORY_TIME_OUTPUT` | `true` | Create unique output subdirectory with date and launch time |
### Example Hosttrace Instrumentation
#### Binary Rewrite
Rewrite the text section of an executable with instrumentation:
Rewrite the text section of an executable or library with instrumentation:
```console
```shell
hosttrace -o app.inst -- /path/to/app
```
In binary rewrite mode, if you also want instrumentation in the linked libraries, you must also rewrite those libraries.
Example of rewriting the functions starting with `"hip"` with instrumentation in the amdhip64 library:
```console
```shell
mkdir -p ./lib
hosttrace -R '^hip' -o ./lib/libamdhip64.so.4 -- /opt/rocm/lib/libamdhip64.so.4
export LD_LIBRARY_PATH=${PWD}/lib:${LD_LIBRARY_PATH}
@@ -76,9 +92,32 @@ export LD_LIBRARY_PATH=${PWD}/lib:${LD_LIBRARY_PATH}
Once you have rewritten your executable and/or libraries with instrumentation, you can just run the (instrumented) executable
or exectuable which loads the instrumented libraries normally, e.g.:
```console
```shell
./app.inst
```
If you want to re-define certain settings to new default in a binary rewrite, use the `--env` option. This `hosttrace` option
will set the environment variable to the given value but will not override it. E.g. the default value of `HOSTTRACE_BUFFER_SIZE_KB`
is 1024000 KB (1 GiB):
```shell
# buffer size defaults to 1024000
hosttrace -o app.inst -- /path/to/app
./app.inst
```
Passing `--env HOSTTRACE_BUFFER_SIZE_KB=5120000` will change the default value in `app.inst` to 5120000 KiB (5 GiB):
```shell
# defaults to 5 GiB buffer size
hosttrace -o app.inst --env HOSTTRACE_BUFFER_SIZE_KB=5120000 -- /path/to/app
./app.inst
```
```shell
# override default 5 GiB buffer size to 200 MB
export HOSTTRACE_BUFFER_SIZE_KB=200000
./app.inst
rocprof --hip-trace --roctx-trace --stats ./app.inst
```
#### Runtime Instrumentation
@@ -86,7 +125,7 @@ rocprof --hip-trace --roctx-trace --stats ./app.inst
Runtime instrumentation will not only instrument the text section of the executable but also the text sections of the
linked libraries. Thus, it may be useful to exclude those libraries via the `-ME` (module exclude) regex option.
```console
```shell
hosttrace -- /path/to/app
hosttrace -ME '^(libhsa-runtime64|libz\\.so)' -- /path/to/app
hosttrace -E 'rocr::atomic|rocr::core|rocr::HSA' -- /path/to/app
@@ -94,8 +133,12 @@ hosttrace -E 'rocr::atomic|rocr::core|rocr::HSA' -- /path/to/app
## Miscellaneous Features and Caveats
- You may need to increase the default perfetto buffer size (1 GB) to capture all the information
- E.g. `export HOSTTRACE_BUFFER_SIZE_KB=10240000`
- You may need to increase the default perfetto buffer size (1 GiB) to capture all the information
- E.g. `export HOSTTRACE_BUFFER_SIZE_KB=10240000` increases the buffer size to 10 GiB
- The hosttrace library has various setting which can be configured via environment variables, you can
configure these settings to custom defaults with the hosttrace command-line tool via the `--env` option
- E.g. to default to a buffer size of 5 GB, use `--env HOSTTRACE_BUFFER_SIZE_KB=5120000`
- This is particularly useful in binary rewrite mode
- Perfetto tooling is enabled by default
- Timemory tooling is disabled by default
- Enabling/disabling one of the aformentioned tools but not specifying enabling/disable the other will assume the inverse of the other's enabled state, e.g.
@@ -122,9 +165,16 @@ variable. The special character sequences `%pid%` and `%rank%` will be replaced
## Merging the traces from rocprof and hosttrace
> NOTE: Using `rocprof` externally is deprecated. The current version has built-in support for
> recording the GPU activity and HIP API calls. If you want to use an external rocprof, either
> configure CMake with `-DHOSTTRACE_USE_ROCTRACER=OFF` or explicitly set `TIMEMORY_ROCTRACER_ENABLED=OFF` in the
> environment.
Use the `hosttrace-merge.jl` Julia script to merge rocprof and perfetto traces.
```console
```shell
export TIMEMORY_ROCTRACER_ENABLED=OFF
rocprof --hip-trace --roctx-trace --stats ./app.inst
hosttrace-merge.jl results.json hosttrace-app.inst-output/2021-09-02_01.03_PM/*.proto
```
@@ -132,7 +182,7 @@ hosttrace-merge.jl results.json hosttrace-app.inst-output/2021-09-02_01.03_PM/*.
In a separate window run:
```console
```shell
pkill traced
traced --background
perfetto --out ./htrace.out --txt -c ${HOSTTRACE_ROOT}/share/roctrace.cfg
@@ -140,12 +190,12 @@ perfetto --out ./htrace.out --txt -c ${HOSTTRACE_ROOT}/share/roctrace.cfg
then in the window running the application, configure the hosttrace instrumentation to use the system backend:
```console
```shell
export HOSTTRACE_BACKEND_SYSTEM=1
```
for the merge use the `htrace.out`:
```console
```shell
hosttrace-merge.jl results.json htrace.out
```
cmake/Compilers.cmake
+7 -1
View File
@@ -535,11 +535,17 @@ foreach(LANG C CXX CUDA)
SET_COMPILER_VAR( CLANG 1)
SET_COMPILER_VAR( APPLE_CLANG 1)
elseif(CMAKE_${LANG}_COMPILER_ID MATCHES "Clang")
# Clang/LLVM compiler
SET_COMPILER_VAR( CLANG 1)
# HIP Clang compiler
if(CMAKE_${LANG}_COMPILER MATCHES "hipcc")
SET_COMPILER_VAR( HIPCC 1)
endif()
elseif(CMAKE_${LANG}_COMPILER_ID MATCHES "PGI")
# PGI compiler
@@ -585,7 +591,7 @@ foreach(LANG C CXX CUDA)
endif()
# set other to no
foreach(TYPE GNU INTEL INTEL_ICC INTEL_ICPC APPLE_CLANG CLANG PGI XLC HP_ACC MIPS MSVC NVIDIA)
foreach(TYPE GNU INTEL INTEL_ICC INTEL_ICPC APPLE_CLANG CLANG PGI XLC HP_ACC MIPS MSVC NVIDIA HIPCC)
if(NOT DEFINED CMAKE_${LANG}_COMPILER_IS_${TYPE})
SET_COMPILER_VAR(${TYPE} 0)
endif()
cmake/Findroctracer.cmake
+104
View File
@@ -0,0 +1,104 @@
# Distributed under the OSI-approved BSD 3-Clause License. See accompanying
# file Copyright.txt or https://cmake.org/licensing for details.
include(FindPackageHandleStandardArgs)
#----------------------------------------------------------------------------------------#
set(_ROCM_PATHS $ENV{ROCM_HOME} /opt/rocm /opt/rocm/roctracer)
#----------------------------------------------------------------------------------------#
find_path(roctracer_ROOT_DIR
NAMES include/roctracer.h
HINTS ${_ROCM_PATHS}
PATHS ${_ROCM_PATHS}
PATH_SUFFIXES roctracer)
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)
mark_as_advanced(roctracer_INCLUDE_DIR)
#----------------------------------------------------------------------------------------#
find_library(roctracer_LIBRARY
NAMES roctracer64 roctracer
HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_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}
PATH_SUFFIXES lib lib64)
find_library(roctracer_kfdwrapper_LIBRARY
NAMES kfdwrapper64 kfdwrapper
HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
PATH_SUFFIXES lib lib64)
find_library(roctracer_hsakmt_LIBRARY
NAMES hsakmt
HINTS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
PATHS ${roctracer_ROOT_DIR} ${_ROCM_PATHS}
PATH_SUFFIXES lib lib64)
if(roctracer_LIBRARY)
get_filename_component(roctracer_LIBRARY_DIR "${roctracer_LIBRARY}" PATH CACHE)
endif()
mark_as_advanced(roctracer_LIBRARY roctracer_roctx_LIBRARY)
#----------------------------------------------------------------------------------------#
find_package_handle_standard_args(roctracer DEFAULT_MSG
roctracer_ROOT_DIR
roctracer_INCLUDE_DIR
roctracer_LIBRARY
roctracer_roctx_LIBRARY)
#------------------------------------------------------------------------------#
if(roctracer_FOUND)
add_library(roctracer::roctracer INTERFACE IMPORTED)
add_library(roctracer::roctx INTERFACE IMPORTED)
set(roctracer_INCLUDE_DIRS ${roctracer_INCLUDE_DIR})
set(roctracer_LIBRARIES ${roctracer_LIBRARY} ${roctracer_roctx_LIBRARY})
set(roctracer_LIBRARY_DIRS ${roctracer_LIBRARY_DIR})
target_include_directories(roctracer::roctracer INTERFACE ${roctracer_INCLUDE_DIR})
target_include_directories(roctracer::roctx INTERFACE ${roctracer_INCLUDE_DIR})
target_link_libraries(roctracer::roctracer INTERFACE ${roctracer_LIBRARY})
target_link_libraries(roctracer::roctx INTERFACE ${roctracer_roctx_LIBRARY})
if(roctracer_kfdwrapper_LIBRARY)
list(APPEND roctracer_LIBRARIES ${roctracer_kfdwrapper_LIBRARY})
target_compile_definitions(roctracer::roctracer INTERFACE HOSTTRACE_ROCTRACER_LIBKFDWRAPPER=\"${roctracer_kfdwrapper_LIBRARY}\")
target_link_libraries(roctracer::roctracer INTERFACE ${roctracer_kfdwrapper_LIBRARY})
target_link_libraries(roctracer::roctx INTERFACE ${roctracer_kfdwrapper_LIBRARY})
endif()
if(roctracer_hsakmt_LIBRARY)
list(APPEND roctracer_LIBRARIES ${roctracer_hsakmt_LIBRARY})
target_link_libraries(roctracer::roctracer INTERFACE ${roctracer_hsakmt_LIBRARY})
target_link_libraries(roctracer::roctx INTERFACE ${roctracer_hsakmt_LIBRARY})
endif()
endif()
#------------------------------------------------------------------------------#
unset(_ROCM_PATHS)
#------------------------------------------------------------------------------#
cmake/MacroUtilities.cmake
+48
View File
@@ -1522,5 +1522,53 @@ FUNCTION(PRINT_FEATURES)
print_disabled_interfaces()
ENDFUNCTION()
#----------------------------------------------------------------------------------------#
# this function is provided to easily select which files use alternative compiler:
#
# GLOBAL --> all files
# TARGET --> all files in a target
# SOURCE --> specific source files
# DIRECTORY --> all files in directory
# PROJECT --> all files/targets in a project/subproject
#
FUNCTION(hosttrace_custom_compilation)
CMAKE_PARSE_ARGUMENTS(COMP "GLOBAL;PROJECT" "COMPILER" "DIRECTORY;TARGET;SOURCE" ${ARGN})
# find hosttrace_launch_compiler
FIND_PROGRAM(HOSTTRACE_COMPILE_LAUNCHER
NAMES hosttrace_launch_compiler
HINTS ${PROJECT_SOURCE_DIR} ${CMAKE_SOURCE_DIR}
PATHS ${PROJECT_SOURCE_DIR} ${CMAKE_SOURCE_DIR}
PATH_SUFFIXES scripts bin)
IF(NOT COMP_COMPILER)
MESSAGE(FATAL_ERROR "hosttrace_custom_compilation not provided COMPILER argument")
ENDIF()
IF(NOT HOSTTRACE_COMPILE_LAUNCHER)
MESSAGE(FATAL_ERROR "hosttrace could not find 'hosttrace_launch_compiler'. Please set '-DHOSTTRACE_COMPILE_LAUNCHER=/path/to/launcher'")
ENDIF()
IF(COMP_GLOBAL)
# if global, don't bother setting others
SET_PROPERTY(GLOBAL PROPERTY RULE_LAUNCH_COMPILE "${HOSTTRACE_COMPILE_LAUNCHER} ${COMP_COMPILER} ${CMAKE_CXX_COMPILER}")
SET_PROPERTY(GLOBAL PROPERTY RULE_LAUNCH_LINK "${HOSTTRACE_COMPILE_LAUNCHER} ${COMP_COMPILER} ${CMAKE_CXX_COMPILER}")
ELSE()
FOREACH(_TYPE PROJECT DIRECTORY TARGET SOURCE)
# make project/subproject scoping easy, e.g. hosttrace_custom_compilation(PROJECT) after project(...)
IF("${_TYPE}" STREQUAL "PROJECT" AND COMP_${_TYPE})
LIST(APPEND COMP_DIRECTORY ${PROJECT_SOURCE_DIR})
UNSET(COMP_${_TYPE})
ENDIF()
# set the properties if defined
IF(COMP_${_TYPE})
FOREACH(_VAL ${COMP_${_TYPE}})
SET_PROPERTY(${_TYPE} ${_VAL} PROPERTY RULE_LAUNCH_COMPILE "${HOSTTRACE_COMPILE_LAUNCHER} ${COMP_COMPILER} ${CMAKE_CXX_COMPILER}")
SET_PROPERTY(${_TYPE} ${_VAL} PROPERTY RULE_LAUNCH_LINK "${HOSTTRACE_COMPILE_LAUNCHER} ${COMP_COMPILER} ${CMAKE_CXX_COMPILER}")
ENDFOREACH()
ENDIF()
ENDFOREACH()
ENDIF()
ENDFUNCTION()
cmake_policy(POP)
cmake/Packages.cmake
+47 -10
View File
@@ -12,7 +12,9 @@ add_interface_library(hosttrace-headers
add_interface_library(hosttrace-threading
"Enables multithreading support")
add_interface_library(hosttrace-dyninst
"Provides flags and libraries for Dyninst (dynamic instrumentation")
"Provides flags and libraries for Dyninst (dynamic instrumentation)")
add_interface_library(hosttrace-roctracer
"Provides flags and libraries for roctracer")
# include threading because of rooflines
target_link_libraries(hosttrace-headers INTERFACE hosttrace-threading)
@@ -41,13 +43,28 @@ if(pthread_LIBRARY AND NOT WIN32)
endif()
#----------------------------------------------------------------------------------------#
#
# roctracer
#
#----------------------------------------------------------------------------------------#
if(HOSTTRACE_USE_ROCTRACER)
find_package(roctracer ${hosttrace_FIND_QUIETLY} REQUIRED)
find_package(hip ${hosttrace_FIND_QUIETLY} REQUIRED)
target_compile_definitions(hosttrace-roctracer INTERFACE HOSTTRACE_USE_ROCTRACER)
target_link_libraries(hosttrace-roctracer INTERFACE hip::host roctracer::roctracer)
set(CMAKE_INSTALL_RPATH "${CMAKE_INSTALL_RPATH}:${roctracer_LIBRARY_DIRS}")
endif()
#----------------------------------------------------------------------------------------#
#
# Dyninst
#
#----------------------------------------------------------------------------------------#
find_package(Dyninst ${hosttrace_FIND_QUIETLY} REQUIRED)
find_package(Dyninst ${hosttrace_FIND_QUIETLY} REQUIRED
COMPONENTS dyninstAPI parseAPI instructionAPI symtabAPI)
set(_BOOST_COMPONENTS atomic system thread date_time)
set(hosttrace_BOOST_COMPONENTS "${_BOOST_COMPONENTS}" CACHE STRING
"Boost components used by Dyninst in hosttrace")
@@ -55,7 +72,6 @@ set(Boost_NO_BOOST_CMAKE ON)
find_package(Boost QUIET REQUIRED
COMPONENTS ${hosttrace_BOOST_COMPONENTS})
# some installs of dyninst don't set this properly
if(EXISTS "${DYNINST_INCLUDE_DIR}" AND NOT DYNINST_HEADER_DIR)
get_filename_component(DYNINST_HEADER_DIR "${DYNINST_INCLUDE_DIR}" REALPATH CACHE)
@@ -97,9 +113,12 @@ endif()
add_rpath(${DYNINST_LIBRARIES} ${Boost_LIBRARIES})
target_link_libraries(hosttrace-dyninst INTERFACE
${DYNINST_LIBRARIES} ${Boost_LIBRARIES})
foreach(_TARG Dyninst::dyninst Boost::headers Boost::atomic
Boost::system Boost::thread Boost::date_time)
if(TARGET ${_TARG})
foreach(_TARG dyninst dyninstAPI instructionAPI symtabAPI parseAPI headers atomic system thread date_time TBB)
if(TARGET Dyninst::${_TARG})
target_link_libraries(hosttrace-dyninst INTERFACE Dyninst::${_TARG})
elseif(TARGET Boost::${_TARG})
target_link_libraries(hosttrace-dyninst INTERFACE Boost::${_TARG})
elseif(TARGET ${_TARG})
target_link_libraries(hosttrace-dyninst INTERFACE ${_TARG})
endif()
endforeach()
@@ -129,6 +148,22 @@ checkout_git_submodule(
REPO_BRANCH v17.0
TEST_FILE sdk/perfetto.cc)
#----------------------------------------------------------------------------------------#
#
# ELFIO
#
#----------------------------------------------------------------------------------------#
if(HOSTTRACE_BUILD_DEVICETRACE)
checkout_git_submodule(
RELATIVE_PATH external/elfio
WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}
REPO_URL https://github.com/jrmadsen/ELFIO.git
REPO_BRANCH set-offset-support)
add_subdirectory(external/elfio)
endif()
#----------------------------------------------------------------------------------------#
#
# Clang Tidy
@@ -165,7 +200,7 @@ endmacro()
#
#------------------------------------------------------------------------------#
find_program(CLANG_FORMAT_EXE
find_program(HOSTTRACE_CLANG_FORMAT_EXE
NAMES
clang-format-12
clang-format-11
@@ -173,7 +208,7 @@ find_program(CLANG_FORMAT_EXE
clang-format-9
clang-format)
if(CLANG_FORMAT_EXE)
if(HOSTTRACE_CLANG_FORMAT_EXE)
file(GLOB sources
${PROJECT_SOURCE_DIR}/src/*.cpp)
file(GLOB headers
@@ -182,8 +217,8 @@ if(CLANG_FORMAT_EXE)
${PROJECT_SOURCE_DIR}/examples/*.cpp
${PROJECT_SOURCE_DIR}/examples/*.hpp)
add_custom_target(format
${CLANG_FORMAT_EXE} -i ${sources} ${headers} ${examples}
COMMENT "Running ${CLANG_FORMAT_EXE}...")
${HOSTTRACE_CLANG_FORMAT_EXE} -i ${sources} ${headers} ${examples}
COMMENT "Running ${HOSTTRACE_CLANG_FORMAT_EXE}...")
else()
message(AUTHOR_WARNING "clang-format could not be found. format build target not available.")
endif()
@@ -195,6 +230,8 @@ endif()
set(TIMEMORY_INSTALL_HEADERS OFF CACHE BOOL "Disable timemory header install")
set(TIMEMORY_INSTALL_CONFIG OFF CACHE BOOL "Disable timemory cmake configuration install")
set(TIMEMORY_INSTALL_ALL OFF CACHE BOOL "Disable install target depending on all target")
set(TIMEMORY_BUILD_C OFF CACHE BOOL "Disable timemory C library")
set(TIMEMORY_BUILD_FORTRAN OFF CACHE BOOL "Disable timemory Fortran library")
set(TIMEMORY_BUILD_TOOLS OFF CACHE BOOL "Ensure timem executable is built")
set(TIMEMORY_BUILD_EXCLUDE_FROM_ALL ON CACHE BOOL "Set timemory to only build dependencies")
set(TIMEMORY_QUIET_CONFIG ON CACHE BOOL "Make timemory configuration quieter")
examples/parallel-overhead/CMakeLists.txt
+10 -1
View File
@@ -1,5 +1,14 @@
cmake_minimum_required(VERSION 3.13 FATAL_ERROR)
project(hosttrace-parallel-overhead
LANGUAGES CXX)
set(CMAKE_BUILD_TYPE "Release")
set(CMAKE_RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}")
find_package(Threads REQUIRED)
add_executable(parallel-overhead parallel-overhead.cpp)
target_link_libraries(parallel-overhead Threads::Threads)
if(NOT CMAKE_PROJECT_NAME STREQUAL PROJECT_NAME)
set_target_properties(transpose PROPERTIES
RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
endif()
examples/transpose/CMakeLists.txt
+21 -84
View File
@@ -1,106 +1,43 @@
cmake_minimum_required(VERSION 3.13 FATAL_ERROR)
# uses make to avoid CMAKE_CXX_COMPILER issues
project(transpose
project(hosttrace-transpose
LANGUAGES CXX)
find_program(MAKE_EXECUTABLE NAMES make gmake)
find_program(HIPCC_EXECUTABLE NAMES hipcc)
if(NOT MAKE_EXECUTABLE)
message(AUTHOR_WARNING "make/gmake could not be found. Cannot build transpose target")
return()
endif()
if(NOT HIPCC_EXECUTABLE)
message(AUTHOR_WARNING "hipcc could not be found. Cannot build transpose target")
return()
endif()
string(TOUPPER "${CMAKE_BUILD_TYPE}" BUILD_TYPE)
set(transpose_CXX_FLAGS "-W -Wall ${CMAKE_CXX_FLAGS_${BUILD_TYPE}} ${CMAKE_CXX_FLAGS}")
set(transpose_LINK_FLAGS "")
if(CMAKE_CXX_COMPILER_IS_CLANG AND TARGET hosttrace::hosttrace-compile-options)
set(transpose_CXX_FLAGS "${transpose_CXX_FLAGS} ${hosttrace_CXX_FLAGS}")
get_target_property(COMPILE_OPTS hosttrace::hosttrace-compile-options INTERFACE_COMPILE_OPTIONS)
set(transpose_CXX_FLAGS "${transpose_CXX_FLAGS} ${COMPILE_OPTS}")
endif()
option(TRANSPOSE_USE_MPI "Enable MPI support in transpose exe" ${TIMEMORY_USE_MPI})
if(TRANSPOSE_USE_MPI)
find_package(MPI REQUIRED)
endif()
if(TARGET MPI::MPI_C)
set(transpose_CXX_FLAGS "${transpose_CXX_FLAGS} -DUSE_MPI")
get_target_property(COMPILE_OPTS MPI::MPI_C INTERFACE_COMPILE_OPTIONS)
foreach(_COPT ${COMPILE_OPTS})
if(NOT "${_COPT}" MATCHES "COMPILE_LANG")
set(transpose_CXX_FLAGS "${transpose_CXX_FLAGS} ${_COPT}")
endif()
endforeach()
get_target_property(INCLUDE_DIRS MPI::MPI_C INTERFACE_INCLUDE_DIRECTORIES)
foreach(_IDIR ${INCLUDE_DIRS})
set(transpose_CXX_FLAGS "${transpose_CXX_FLAGS} -I${_IDIR}")
endforeach()
if(MPI_C_LINK_FLAGS)
set(transpose_LINK_FLAGS "${transpose_LINK_FLAGS} ${MPI_C_LINK_FLAGS}")
endif()
set(_LINK_LIBS "")
foreach(_LIB ${MPI_C_LIB_NAMES})
string(APPEND _LINK_LIBS "-l${_LIB} ")
endforeach()
foreach(_IDIR ${INCLUDE_DIRS} ${MPI_mpich_LIBRARY} ${MPI_mpi_LIBRARY} ${MPI_LIBRARY_DIRS})
get_filename_component(_LIBDIR "${_IDIR}" DIRECTORY)
if(EXISTS "${_IDIR}/libmpi${CMAKE_SHARED_LIBRARY_SUFFIX}")
set(transpose_LINK_FLAGS "${transpose_LINK_FLAGS} -L${_IDIR} ${_LINK_LIBS}")
endif()
if(EXISTS "${_LIBDIR}/libmpi${CMAKE_SHARED_LIBRARY_SUFFIX}")
set(transpose_LINK_FLAGS "${transpose_LINK_FLAGS} -L${_LIBDIR} ${_LINK_LIBS}")
endif()
foreach(_LDIR lib lib64)
set(_LIBDIR_SAVE "${_LIBDIR}")
if(NOT EXISTS "${_LIBDIR}/${_LDIR}")
get_filename_component(_LIBDIR "${_LIBDIR}" DIRECTORY)
endif()
if(EXISTS "${_LIBDIR}/${_LDIR}")
set(transpose_LINK_FLAGS "${transpose_LINK_FLAGS} -L${_LIBDIR}/${_LDIR} ${_LINK_LIBS}")
endif()
set(_LIBDIR "${_LIBDIR_SAVE}")
endforeach()
endforeach()
add_executable(transpose transpose.cpp)
if(CMAKE_CXX_COMPILER_IS_CLANG AND TARGET hosttrace::hosttrace-compile-options)
target_link_libraries(transpose PRIVATE hosttrace::hosttrace-compile-options)
else()
target_compile_options(transpose PRIVATE -W -Wall)
endif()
# remove generator expressions
string(REPLACE "$<" "" transpose_CXX_FLAGS "${transpose_CXX_FLAGS}")
string(REPLACE ">:" ":" transpose_CXX_FLAGS "${transpose_CXX_FLAGS}")
string(REPLACE ">" "" transpose_CXX_FLAGS "${transpose_CXX_FLAGS}")
string(REPLACE "COMPILE_LANGUAGE:CXX:" "" transpose_CXX_FLAGS "${transpose_CXX_FLAGS}")
# remove double spaces
string(REPLACE " " " " transpose_CXX_FLAGS "${transpose_CXX_FLAGS}")
# convert to list
string(REPLACE " " ";" transpose_CXX_FLAGS "${transpose_CXX_FLAGS}")
# remove duplicate args
list(REMOVE_DUPLICATES transpose_CXX_FLAGS)
# make single string
string(REPLACE ";" " " transpose_CXX_FLAGS "${transpose_CXX_FLAGS}")
if(TRANSPOSE_USE_MPI)
target_compile_definitions(transpose PRIVATE USE_MPI)
target_link_libraries(transpose PRIVATE MPI::MPI_C)
endif()
# configure files in build directory
configure_file(${PROJECT_SOURCE_DIR}/transpose.cpp
${CMAKE_BINARY_DIR}/transpose/transpose.cpp COPYONLY)
if(NOT CMAKE_PROJECT_NAME STREQUAL PROJECT_NAME)
set_target_properties(transpose PROPERTIES
RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
configure_file(${PROJECT_SOURCE_DIR}/Makefile.in
${CMAKE_BINARY_DIR}/transpose/Makefile @ONLY)
# touch the file when cmake re-runs to ensure rebuilt
execute_process(COMMAND ${CMAKE_COMMAND} -E touch ${CMAKE_BINARY_DIR}/transpose/transpose.cpp)
# target for adding a dependency to libpytimemory target
add_custom_target(transpose ALL
COMMAND ${MAKE_EXECUTABLE}
WORKING_DIRECTORY ${CMAKE_BINARY_DIR}/transpose
BYPRODUCTS ${CMAKE_BINARY_DIR}/transpose/transpose
SOURCES ${CMAKE_BINARY_DIR}/transpose/transpose.cpp)
if(NOT CMAKE_CXX_COMPILER_IS_HIPCC AND HIPCC_EXECUTABLE)
# defined in MacroUtilities.cmake
hosttrace_custom_compilation(
COMPILER ${HIPCC_EXECUTABLE}
TARGET transpose)
endif()
endif()
examples/transpose/Makefile
-40
View File
@@ -1,40 +0,0 @@
CXXFLAGS += -std=c++17 -O3 -Wno-unused-result -g
LINK_FLAGS :=
all: transpose
debug: CXXFLAGS += -ggdb -g
debug: transpose
transpose: transpose.cpp
hipcc ${CXXFLAGS} -o $@ $^ ${LINK_FLAGS}
.PHONY: clean
run:
./transpose
run.verbose:
AMD_LOG_LEVEL=3, AMD_LOG_MASK=128 ./transpose
rd:
rocgdb --args ./transpose
prof.trace:
rocprof --hip-trace --stats ./transpose
prof:
rocprof --obj-tracking on -i profile.txt ./transpose
prof.full:
rocprof --obj-tracking on -i profile.txt -m custom_metric.xml ./transpose
callgrind:
valgrind --tool=callgrind ./transpose
callgrind.ann:
callgrind_annotate --inclusive=yes --tree=both callgrind.out.* > callgrind.out
clean:
rm -f transpose *.o results.*
examples/transpose/Makefile.in
-39
View File
@@ -1,39 +0,0 @@
CXXFLAGS += -std=c++@CMAKE_CXX_STANDARD@ @transpose_CXX_FLAGS@
LINKFLAGS := @transpose_LINK_FLAGS@
all: transpose
debug: CXXFLAGS += -ggdb -g
debug: transpose
transpose: transpose.cpp
hipcc ${CXXFLAGS} -o $@ $^ ${LINKFLAGS}
.PHONY: clean
run:
./transpose
run.verbose:
AMD_LOG_LEVEL=3, AMD_LOG_MASK=128 ./transpose
rd:
rocgdb --args ./transpose
prof.trace:
rocprof --hip-trace --stats ./transpose
prof:
rocprof --obj-tracking on -i profile.txt ./transpose
prof.full:
rocprof --obj-tracking on -i profile.txt -m custom_metric.xml ./transpose
callgrind:
valgrind --tool=callgrind ./transpose
callgrind.ann:
callgrind_annotate --inclusive=yes --tree=both callgrind.out.* > callgrind.out
clean:
rm -f transpose *.o results.*
examples/transpose/transpose.cpp
+4 -3
View File
@@ -123,14 +123,15 @@ run(int argc, char** argv)
for(size_t i = 0; i < times; i++)
{
hipLaunchKernelGGL(transpose_a, grid, block, 0, 0, in, out, M, N);
check_hip_error();
HIP_API_CALL(hipDeviceSynchronize());
}
check_hip_error();
HIP_API_CALL(hipDeviceSynchronize());
t2 = std::chrono::high_resolution_clock::now();
double time =
std::chrono::duration_cast<std::chrono::duration<double>>(t2 - t1).count();
float GB = (float) size * times * 2 / (1 << 30);
std::cout << "The average performance of transpose is " << GB / time << " GBytes/sec"
std::cout << "Runtime of transpose is " << time << " sec\n"
<< "The average performance of transpose is " << GB / time << " GBytes/sec"
<< std::endl;
int* out_matrix = (int*) malloc(size);
external/elfio
Vendored Submodule
+1
Submodule external/elfio added at 637221d267
external/timemory
Vendored
+1 -1
Submodule external/timemory updated: aa4a0ed7b2...cc0cfd6101
include/hosttrace.hpp
-6
View File
@@ -396,9 +396,6 @@ struct module_function
std::pair<address_t, address_t> _range{};
if(proc->getAddressRange(_range.first, _range.second))
address_range = _range.second - _range.first;
auto _instructions = proc->findPoint(BPatch_locInstruction);
if(_instructions)
instr_count = _instructions->size();
}
friend bool operator<(const module_function& lhs, const module_function& rhs)
@@ -418,7 +415,6 @@ struct module_function
std::stringstream ss;
ss << std::setw(14) << "AddressRange"
<< " " << std::setw(14) << "InstrCount"
<< " " << std::setw(w0 + 8) << std::left << "Module"
<< " " << std::setw(w1 + 8) << std::left << "Function"
<< " " << std::setw(w2 + 8) << std::left << "FunctionSignature"
@@ -442,7 +438,6 @@ struct module_function
// clang-format off
ss << std::setw(14) << rhs.address_range << " "
<< std::setw(14) << rhs.instr_count << " "
<< std::setw(w0 + 8) << std::left << _get_str(rhs.module) << " "
<< std::setw(w1 + 8) << std::left << _get_str(rhs.function) << " "
<< std::setw(w2 + 8) << std::left << _get_str(rhs.signature.get());
@@ -453,7 +448,6 @@ struct module_function
}
size_t address_range = 0;
size_t instr_count = 0;
string_t module = {};
string_t function = {};
function_signature signature;
include/library.hpp
+67 -6
View File
@@ -4,12 +4,14 @@
#if !defined(TIMEMORY_USE_PERFETTO)
# include <perfetto.h>
# define PERFETTO_CATEGORIES \
perfetto::Category("hosttrace").SetDescription("Function trace")
perfetto::Category("host").SetDescription("Host-side function tracing"), \
perfetto::Category("device").SetDescription("Device-side function tracing")
#else
# define PERFETTO_CATEGORIES \
perfetto::Category("hosttrace").SetDescription("Function trace"), \
perfetto::Category("timemory") \
.SetDescription("Events from the timemory API")
perfetto::Category("host").SetDescription("Host-side function tracing"), \
perfetto::Category("device").SetDescription("Device-side function tracing")
perfetto::Category("timemory")
.SetDescription("Events from the timemory API")
# define TIMEMORY_PERFETTO_CATEGORIES PERFETTO_CATEGORIES
#endif
@@ -42,6 +44,8 @@
#include "timemory/storage.hpp"
#include "timemory/variadic.hpp"
#include "roctracer.hpp"
// forward decl of the API
extern "C"
{
@@ -74,6 +78,11 @@ namespace audit = tim::audit;
namespace comp = tim::component;
namespace quirk = tim::quirk;
namespace threading = tim::threading;
namespace scope = tim::scope;
namespace dmp = tim::dmp;
namespace process = tim::process;
namespace units = tim::units;
namespace trait = tim::trait;
// this is used to wrap fork()
struct fork_gotcha : comp::base<fork_gotcha, void>
@@ -122,12 +131,16 @@ private:
const char* m_prefix = nullptr;
};
using papi_tot_ins = comp::papi_tuple<PAPI_TOT_INS>;
using fork_gotcha_t = comp::gotcha<4, tim::component_tuple<fork_gotcha>, hosttrace>;
using mpi_gotcha_t = comp::gotcha<4, tim::component_tuple<mpi_gotcha>, hosttrace>;
using hosttrace_bundle_t =
tim::lightweight_tuple<comp::wall_clock, comp::peak_rss, comp::cpu_clock,
comp::cpu_util, comp::user_global_bundle, fork_gotcha_t,
mpi_gotcha_t>;
comp::cpu_util, comp::roctracer, papi_tot_ins,
comp::user_global_bundle, fork_gotcha_t, mpi_gotcha_t>;
using hosttrace_thread_bundle_t =
tim::lightweight_tuple<comp::wall_clock, comp::thread_cpu_clock,
comp::thread_cpu_util, papi_tot_ins>;
using bundle_t =
tim::component_bundle<hosttrace, comp::wall_clock*, comp::user_global_bundle*>;
using bundle_allocator_t = tim::data::ring_buffer_allocator<bundle_t>;
@@ -190,6 +203,54 @@ get_state();
std::unique_ptr<hosttrace_bundle_t>&
get_main_bundle();
bool
get_use_perfetto();
bool
get_use_timemory();
//--------------------------------------------------------------------------------------//
template <typename Tp, size_t MaxThreads = 1024>
struct hosttrace_thread_data
{
static constexpr size_t max_supported_threads = MaxThreads;
using instance_array_t = std::array<std::unique_ptr<Tp>, max_supported_threads>;
template <typename... Args>
static void construct(Args&&...);
static std::unique_ptr<Tp>& instance();
static instance_array_t& instances();
};
template <typename Tp, size_t MaxThreads>
template <typename... Args>
void
hosttrace_thread_data<Tp, MaxThreads>::construct(Args&&... _args)
{
static thread_local bool _v = [&_args...]() {
instances().at(threading::get_id()) =
std::make_unique<Tp>(std::forward<Args>(_args)...);
return true;
}();
(void) _v;
}
template <typename Tp, size_t MaxThreads>
std::unique_ptr<Tp>&
hosttrace_thread_data<Tp, MaxThreads>::instance()
{
return instances().at(threading::get_id());
}
template <typename Tp, size_t MaxThreads>
typename hosttrace_thread_data<Tp, MaxThreads>::instance_array_t&
hosttrace_thread_data<Tp, MaxThreads>::instances()
{
static auto _v = instance_array_t{};
return _v;
}
//--------------------------------------------------------------------------------------//
// there are currently some strange things that happen with vector<bundle_t> so using
include/roctracer.hpp
+59
View File
@@ -0,0 +1,59 @@
#pragma once
#include "timemory/api.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/os.hpp"
#include "timemory/mpl/type_traits.hpp"
#include "timemory/mpl/types.hpp"
TIMEMORY_DECLARE_COMPONENT(roctracer)
#if !defined(HOSTTRACE_USE_ROCTRACER)
TIMEMORY_DEFINE_CONCRETE_TRAIT(is_available, component::roctracer, false_type)
#endif
namespace tim
{
namespace component
{
using roctracer_data = data_tracker<double, roctracer>;
struct roctracer
: base<roctracer, void>
, private policy::instance_tracker<roctracer, false>
{
using value_type = void;
using base_type = base<roctracer, void>;
using tracker_type = policy::instance_tracker<roctracer, false>;
TIMEMORY_DEFAULT_OBJECT(roctracer)
static void setup();
static void tear_down();
static void preinit();
static void global_init() { setup(); }
static void global_finalize() { tear_down(); }
void start();
void stop();
void set_prefix(const char* _v) { m_prefix = _v; }
private:
const char* m_prefix = nullptr;
};
} // namespace component
} // namespace tim
TIMEMORY_SET_COMPONENT_API(component::roctracer_data, project::timemory, category::timing,
os::supports_unix)
TIMEMORY_DEFINE_CONCRETE_TRAIT(is_timing_category, component::roctracer_data, true_type)
TIMEMORY_DEFINE_CONCRETE_TRAIT(uses_timing_units, component::roctracer_data, true_type)
#include "timemory/operations.hpp"
TIMEMORY_DECLARE_EXTERN_COMPONENT(roctracer, false, void)
TIMEMORY_DECLARE_EXTERN_COMPONENT(roctracer_data, true, double)
scripts/gperftools-cpu-profile.sh
Executable
+138
View File
@@ -0,0 +1,138 @@
#!/bin/bash
EXE=$(basename ${1})
DIR=cpu.prof.${EXE}
mkdir -p ${DIR}
# gperf settings
: ${N:=0}
: ${GPERF_PROFILE:=""}
: ${GPERF_PROFILE_BASE:=${DIR}/gperf}
: ${MALLOCSTATS:=1}
: ${CPUPROFILE_FREQUENCY:=250}
: ${CPUPROFILE_REALTIME:=1}
# rendering settings
: ${INTERACTIVE:=0}
: ${IMG_FORMAT:="png"}
#: ${DOT_ARGS:='-Gsize=24,24\! -Gdpi=200'}
: ${DOT_ARGS:=""}
: ${PPROF_ARGS:="--no_strip_temp --functions"}
if [ "$(uname)" = "Darwin" ]; then
if [ "${IMG_FORMAT}" = "jpeg" ]; then
IMG_FORMAT="jpg"
fi
fi
run-verbose()
{
echo "${@}" 1>&2
eval ${@}
}
while [ -z "${GPERF_PROFILE}" ]
do
TEST_FILE=${GPERF_PROFILE_BASE}.${N}
if [ ! -f "${TEST_FILE}" ]; then
GPERF_PROFILE=${TEST_FILE}
fi
N=$((${N}+1))
done
export MALLOCSTATS
export CPUPROFILE_FREQUENCY
export CPUPROFILE_REALTIME
echo -e "\n\t--> Outputting profile to '${GPERF_PROFILE}'...\n"
# remove profile file if unsucessful execution
cleanup-failure() { set +v ; echo "failure"; rm -f ${GPERF_PROFILE}; exit 1; }
trap cleanup-failure SIGHUP SIGINT SIGQUIT SIGILL SIGABRT SIGKILL
ADD_LIBS()
{
for i in $@
do
if [ -z "${ADD_LIB_LIST}" ]; then
ADD_LIB_LIST="--add_lib=${i}"
else
ADD_LIB_LIST="${ADD_LIB_LIST} --add_lib=${i}"
fi
done
}
ADD_PRELOAD()
{
for i in $@
do
if [ -z "${LIBS}" ]; then
LIBS=${i}
else
LIBS="${LIBS}:${i}"
fi
done
}
# 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
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
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!"
fi
else
echo -e "profile file \"${GPERF_PROFILE}\" not found!"
ls -la
fi
scripts/hosttrace_launch_compiler
Executable
+75
View File
@@ -0,0 +1,75 @@
#!/bin/bash -e
#
# This script allows CMAKE_CXX_COMPILER to be a standard
# C++ compiler and hosttrace sets RULE_LAUNCH_COMPILE and
# RULE_LAUNCH_LINK in CMake so that all compiler and link
# commands are prefixed with this script followed by the
# C++ compiler. Thus if $1 == $2 then we know the command
# was intended for the C++ compiler and we discard both
# $1 and $2 and redirect the command to linker.
# If $1 != $2 then we know that the command was not intended
# for the C++ compiler and we just discard $1 and launch
# the original command. Examples of when $2 will not equal
# $1 are 'ar', 'cmake', etc. during the linking phase
#
# emit a message about the underlying command executed
: ${DEBUG:=0}
: ${HOSTTRACE_DEBUG_LAUNCH_COMPILER:=${DEBUG}}
debug-message()
{
if [ "${HOSTTRACE_DEBUG_LAUNCH_COMPILER}" -ne 0 ]; then
echo -e "##### $(basename ${BASH_SOURCE[0]}) executing: \"$@\"... #####"
fi
}
# if hosttrace compiler is not passed, someone is probably trying to invoke it directly
if [ -z "${1}" ]; then
echo -e "\n${BASH_SOURCE[0]} was invoked without the hosttrace compiler as the first argument."
echo "This script is not indended to be directly invoked by any mechanism other"
echo -e "than through a RULE_LAUNCH_COMPILE or RULE_LAUNCH_LINK property set in CMake.\n"
exit 1
fi
# if hosttrace compiler is not passed, someone is probably trying to invoke it directly
if [ -z "${2}" ]; then
echo -e "\n${BASH_SOURCE[0]} was invoked without the C++ compiler as the second argument."
echo "This script is not indended to be directly invoked by any mechanism other"
echo -e "than through a RULE_LAUNCH_COMPILE or RULE_LAUNCH_LINK property set in CMake.\n"
exit 1
fi
# if there aren't two args, this isn't necessarily invalid, just a bit strange
if [ -z "${3}" ]; then exit 0; fi
# store the hosttrace compiler
HOSTTRACE_COMPILER=${1}
# remove the hosttrace compiler from the arguments
shift
# store the expected C++ compiler
CXX_COMPILER=${1}
# remove the expected C++ compiler from the arguments
shift
if [[ "${CXX_COMPILER}" != "${1}" ]]; then
debug-message $@
# the command does not depend on hosttrace so just execute the command w/o re-directing to ${HOSTTRACE_COMPILER}
eval $@
else
# the executable is the C++ compiler, so we need to re-direct to ${HOSTTRACE_COMPILER}
if [ ! -f "${HOSTTRACE_COMPILER}" ]; then
echo -e "\nError: the compiler redirect for hosttrace was not found at ${HOSTTRACE_COMPILER}\n"
exit 1
fi
# discard the compiler from the command
shift
debug-message ${HOSTTRACE_COMPILER} $@
# execute ${HOSTTRACE_COMPILER} (again, usually nvcc_wrapper)
${HOSTTRACE_COMPILER} $@
fi
src/hosttrace-details.cpp
+1 -1
View File
@@ -272,7 +272,7 @@ find_function(image_t* app_image, const std::string& _name, strset_t _extra)
}
if(!_func)
verbprintf(0, "timemory-run: Unable to find function %s\n", _name.c_str());
verbprintf(2, "hosttrace: Unable to find function %s\n", _name.c_str());
return _func;
}
src/hosttrace.cpp
+17 -21
View File
@@ -30,10 +30,11 @@
static bool is_driver = false;
static bool allow_overlapping = false;
static bool instr_dynamic_callsites = false;
static size_t batch_size = 50;
static strset_t extra_libs = {};
static size_t min_address_range = (1 << 9); // 512
static size_t min_loop_address_range = (1 << 6); // 64
static size_t min_loop_address_range = (1 << 9); // 512
static std::vector<std::pair<uint64_t, string_t>> hash_ids = {};
static std::map<string_t, bool> use_stubs = {};
static std::map<string_t, procedure_t*> beg_stubs = {};
@@ -282,6 +283,15 @@ main(int argc, char** argv)
.count(1)
.dtype("size_t")
.action([](parser_t& p) { batch_size = p.get<size_t>("batch-size"); });
parser
.add_argument({ "--dynamic-callsites" },
"Force instrumentation if a function has dynamic callsites (e.g. "
"function pointers)")
.max_count(1)
.dtype("boolean")
.action([](parser_t& p) {
instr_dynamic_callsites = p.get<bool>("dynamic-callsites");
});
parser
.add_argument({ "-r", "--min-address-range" },
"If the address range of a function is less than this value, "
@@ -1361,10 +1371,10 @@ main(int argc, char** argv)
if(cfg)
cfg->getOuterLoops(basic_loop);
// if the function has dynamic callsites and we are in binary rewrite mode,
// force the instrumentation
// if the function has dynamic callsites and user specified instrumenting
// dynamic callsites, force the instrumentation
bool _force_instr = false;
if(cfg && binary_rewrite)
if(cfg && instr_dynamic_callsites)
_force_instr = cfg->containsDynamicCallsites();
auto _address_range = module_function{ mod, itr }.address_range;
@@ -1395,26 +1405,12 @@ main(int argc, char** argv)
(unsigned long) min_loop_address_range);
continue;
}
else if(_address_range >= min_loop_address_range &&
_address_range < min_address_range && _has_loop_entries)
{
verbprintf(
1,
"Enabling function [min-loop-address-range]: %s / %s despite not "
"satisfy minimum loop address range (address range = %lu, minimum "
"= %lu) because it has at least one loop (found: %lu)\n",
name.m_name.c_str(), name.get().c_str(),
(unsigned long) _address_range,
(unsigned long) min_loop_address_range,
(unsigned long) _num_loop_entries);
}
else if(_address_range < min_address_range && _force_instr)
else if(_force_instr)
{
verbprintf(1,
"Enabling function [min-address-range]: %s / %s despite not "
"Enabling function [dynamic-callsite]: %s / %s despite not "
"satisfy minimum address range (address range = %lu, minimum "
"= %lu) because contains dynamic callsites which may not be "
"instrumented in binary rewrite mode\n",
"= %lu) because contains dynamic callsites\n",
name.m_name.c_str(), name.get().c_str(),
(unsigned long) _address_range,
(unsigned long) min_address_range);
src/library.cpp
+122 -36
View File
@@ -15,11 +15,7 @@ get_state()
return _v;
}
//--------------------------------------------------------------------------------------//
namespace
{
auto
bool
get_use_perfetto()
{
// if using timemory, default to perfetto being off
@@ -29,7 +25,7 @@ get_use_perfetto()
return _v;
}
auto
bool
get_use_timemory()
{
// default to opposite of whether perfetto setting
@@ -39,11 +35,32 @@ get_use_timemory()
return _v;
}
//--------------------------------------------------------------------------------------//
namespace
{
size_t&
get_sample_rate()
{
static auto _v = tim::get_env<size_t>("HOSTTRACE_SAMPLE_RATE", 1);
return _v;
}
std::vector<bool>&
get_sample_data()
{
static thread_local auto _v = std::vector<bool>{};
return _v;
}
bool&
get_use_mpi()
{
// this does not enable anything particularly useful when not using timemory
static bool _v = tim::get_env("HOSTTRACE_USE_MPI", false, get_use_timemory());
#if defined(TIMEMORY_USE_MPI)
static bool _v = tim::get_env("HOSTTRACE_USE_MPI", false);
#else
static bool _v = false;
#endif
return _v;
}
@@ -74,7 +91,7 @@ auto
ensure_finalization()
{
HOSTTRACE_DEBUG("[%s]\n", __FUNCTION__);
return tim::scope::destructor{ []() { hosttrace_trace_finalize(); } };
return scope::destructor{ []() { hosttrace_trace_finalize(); } };
}
auto&
@@ -103,10 +120,10 @@ get_perfetto_output_filename()
if(_pos != std::string::npos)
_tmp.replace(_pos, _key.length(), std::to_string(_val()));
};
_replace("%pid%", []() { return tim::process::get_id(); });
_replace("%pid%", []() { return process::get_id(); });
_replace("%rank%", []() { return tim::mpi::rank(); });
// backwards compatibility
_replace("%p", []() { return tim::process::get_id(); });
_replace("%p", []() { return process::get_id(); });
return _tmp;
}
@@ -179,7 +196,7 @@ hosttrace_init_tooling()
if(get_use_perfetto() && get_use_mpi())
tim::settings::time_output() = true;
auto _cmd = tim::read_command_line(tim::process::get_id());
auto _cmd = tim::read_command_line(process::get_id());
auto _exe = (_cmd.empty()) ? "hosttrace" : _cmd.front();
auto _pos = _exe.find_last_of('/');
if(_pos < _exe.length() - 1)
@@ -187,6 +204,10 @@ hosttrace_init_tooling()
tim::timemory_init({ _exe }, "hosttrace-");
if(get_sample_rate() < 1)
get_sample_rate() = 1;
get_sample_data().reserve(512);
if(get_use_timemory())
{
comp::user_global_bundle::global_init();
@@ -215,9 +236,13 @@ hosttrace_init_tooling()
}
// always activate gotcha wrappers
auto& _fork_gotcha = get_main_bundle();
_fork_gotcha->start();
assert(_fork_gotcha->get<mpi_gotcha_t>()->get_is_running());
auto& _main_bundle = get_main_bundle();
_main_bundle->start();
assert(_main_bundle->get<mpi_gotcha_t>()->get_is_running());
#if defined(HOSTTRACE_USE_ROCTRACER)
assert(_main_bundle->get<comp::roctracer>() != nullptr);
assert(_main_bundle->get<comp::roctracer>()->get_is_running());
#endif
perfetto::TracingInitArgs args{};
perfetto::TraceConfig cfg{};
@@ -253,8 +278,18 @@ hosttrace_init_tooling()
(void) get_perfetto_output_filename();
}
static auto _thread_init = [_exe]() {
hosttrace_thread_data<hosttrace_thread_bundle_t>::construct(
TIMEMORY_JOIN("", _exe, "/thread-", threading::get_id()),
quirk::config<quirk::auto_start>{});
static thread_local auto _dtor = scope::destructor{ []() {
hosttrace_thread_data<hosttrace_thread_bundle_t>::instance()->stop();
} };
(void) _dtor;
};
// functors for starting and stopping timemory
static auto _push_timemory = [](const char* name) {
_thread_init();
auto& _data = get_timemory_data();
// this generates a hash for the raw string array
auto _hash = tim::add_hash_id(tim::string_view_t{ name });
@@ -264,6 +299,19 @@ hosttrace_init_tooling()
_bundle->start();
};
static auto _push_perfetto = [](const char* name) {
_thread_init();
TRACE_EVENT_BEGIN("host", perfetto::StaticString(name),
[&](perfetto::EventContext ctx) {
// compile-time check
IF_CONSTEXPR(trait::is_available<papi_tot_ins>::value)
{
ctx.event()->set_thread_instruction_count_absolute(
papi_tot_ins::record().at(0));
}
});
};
static auto _pop_timemory = [](const char* name) {
auto& _data = get_timemory_data();
if(_data.bundles.empty())
@@ -278,25 +326,31 @@ hosttrace_init_tooling()
_data.bundles.pop_back();
};
static auto _pop_perfetto = [](const char*) {
TRACE_EVENT_END("host", [&](perfetto::EventContext ctx) {
IF_CONSTEXPR(trait::is_available<papi_tot_ins>::value)
{
ctx.event()->set_thread_instruction_count_absolute(
papi_tot_ins::record().at(0));
}
});
};
if(get_use_perfetto() && get_use_timemory())
{
// if both are used, then use perfetto overload for calling lambda to launch
// timemory
get_functors().first = [](const char* name) {
TRACE_EVENT_BEGIN("hosttrace", perfetto::StaticString(name),
[&](perfetto::EventContext) { _push_timemory(name); });
_push_perfetto(name);
_push_timemory(name);
};
get_functors().second = [](const char* name) {
TRACE_EVENT_END("hosttrace",
[&](perfetto::EventContext) { _pop_timemory(name); });
_pop_timemory(name);
_pop_perfetto(name);
};
}
else if(get_use_perfetto())
{
get_functors().first = [](const char* name) {
TRACE_EVENT_BEGIN("hosttrace", perfetto::StaticString(name));
};
get_functors().second = [](const char*) { TRACE_EVENT_END("hosttrace"); };
get_functors().first = _push_perfetto;
get_functors().second = _pop_perfetto;
}
else if(get_use_timemory())
{
@@ -304,7 +358,7 @@ hosttrace_init_tooling()
get_functors().second = _pop_timemory;
}
if(tim::dmp::rank() == 0)
if(dmp::rank() == 0)
{
tim::print_env(std::cerr,
[](const std::string& _v) { return _v.find("HOSTTRACE_") == 0; });
@@ -339,7 +393,7 @@ hosttrace_init_tooling()
// ends the tracing session
static auto _ensure_finalization = ensure_finalization();
if(tim::dmp::rank() == 0)
if(dmp::rank() == 0)
puts("");
return true;
}
@@ -366,7 +420,12 @@ extern "C"
HOSTTRACE_DEBUG("[%s] %s\n", __FUNCTION__, name);
}
get_functors().first(name);
static auto _sample_rate = std::max<size_t>(get_sample_rate(), 1);
static thread_local size_t _sample_idx = 0;
auto _enabled = (_sample_idx++ % _sample_rate == 0);
get_sample_data().emplace_back(_enabled);
if(_enabled)
get_functors().first(name);
}
void hosttrace_pop_trace(const char* name)
@@ -374,7 +433,13 @@ extern "C"
if(get_state() == State::Active)
{
HOSTTRACE_DEBUG("[%s] %s\n", __FUNCTION__, name);
get_functors().second(name);
auto& _sample_data = get_sample_data();
if(!_sample_data.empty())
{
if(_sample_data.back())
get_functors().second(name);
_sample_data.pop_back();
}
}
else
{
@@ -396,22 +461,38 @@ extern "C"
HOSTTRACE_DEBUG("[%s]\n", __FUNCTION__);
if(tim::dmp::rank() == 0)
if(dmp::rank() == 0)
puts("");
get_state() = State::Finalized;
// stop the main bundle and report the high-level metrics
if(get_main_bundle())
{
get_main_bundle()->stop();
int64_t _id = (get_use_mpi()) ? tim::dmp::rank() : tim::process::get_id();
int64_t _id = (get_use_mpi()) ? dmp::rank() : process::get_id();
std::stringstream _ss{};
_ss << "[" << __FUNCTION__ << "][" << _id << "] " << *get_main_bundle()
<< "\n";
std::cout << _ss.str();
std::cerr << _ss.str();
get_main_bundle().reset();
}
// print out thread-data if they are not still running
// if they are still running (e.g. thread-pool still alive), the
// thread-specific data will be wrong if try to stop them from
// the main thread.
for(auto& itr : hosttrace_thread_data<hosttrace_thread_bundle_t>::instances())
{
if(itr && itr->get<comp::wall_clock>() &&
!itr->get<comp::wall_clock>()->get_is_running())
{
std::stringstream _ss{};
_ss << *itr << "\n";
std::cerr << _ss.str();
}
}
// ensure that all the MT instances are flushed
for(auto& itr : hosttrace_timemory_data::instances())
{
@@ -443,9 +524,14 @@ extern "C"
return;
}
// Write the trace into a file.
fprintf(stderr, "[%s]> Outputting '%s'. Trace data: %lu bytes...\n",
fprintf(stderr,
"[%s]> Outputting '%s'. Trace data: %lu B (%.2f KB / %.2f MB / %.2f "
"GB)...\n",
__FUNCTION__, get_perfetto_output_filename().c_str(),
(unsigned long) trace_data.size());
(unsigned long) trace_data.size(),
static_cast<double>(trace_data.size()) / units::KB,
static_cast<double>(trace_data.size()) / units::MB,
static_cast<double>(trace_data.size()) / units::GB);
std::ofstream output{};
output.open(get_perfetto_output_filename(), std::ios::out | std::ios::binary);
if(!output)
@@ -472,8 +558,8 @@ extern "C"
(attached) ? "y" : "n");
if(use && !attached)
{
auto& _fork_gotcha = get_main_bundle();
_fork_gotcha->start();
auto& _main_bundle = get_main_bundle();
_main_bundle->start();
tim::set_env("HOSTTRACE_USE_MPI", "ON", 1);
get_use_mpi() = true;
get_state() = State::DelayedInit;
src/roctracer.cpp
+317
View File
@@ -0,0 +1,317 @@
#include "roctracer.hpp"
#include "library.hpp"
#include <roctracer.h>
#include <roctracer_hcc.h>
#include <roctracer_hip.h>
// Macro to check ROC-tracer calls status
#define ROCTRACER_CALL(call) \
do \
{ \
int err = call; \
if(err != 0) \
{ \
std::cerr << roctracer_error_string() << " in: " << #call << std::flush; \
} \
} while(0)
using roctracer_bundle_t = tim::component_tuple<comp::roctracer_data>;
namespace units = tim::units;
namespace
{
static auto&
get_roctracer_kernels()
{
static auto _v = std::unordered_set<uint64_t>{};
return _v;
}
static auto&
get_roctracer_data_map()
{
static auto _v = std::unordered_map<uint64_t, roctracer_bundle_t>{};
return _v;
}
static auto&
get_roctracer_key_map()
{
static auto _v = std::unordered_map<uint64_t, const char*>{};
return _v;
}
using data_type_mutex_t = std::decay_t<decltype(get_roctracer_data_map())>;
// HIP API callback function
void
hip_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg)
{
if(!trait::runtime_enabled<comp::roctracer>::get())
return;
const hip_api_data_t* data = reinterpret_cast<const hip_api_data_t*>(callback_data);
HOSTTRACE_DEBUG("<%-30s id(%u)\tcorrelation_id(%lu) %s>\n",
roctracer_op_string(ACTIVITY_DOMAIN_HIP_API, cid, 0), cid,
data->correlation_id,
(data->phase == ACTIVITY_API_PHASE_ENTER) ? "on-enter" : "on-exit");
if(data->phase == ACTIVITY_API_PHASE_ENTER)
{
switch(cid)
{
case HIP_API_ID___hipPushCallConfiguration:
case HIP_API_ID___hipPopCallConfiguration: break;
case HIP_API_ID_hipLaunchKernel: {
const char* _name =
hipKernelNameRefByPtr(data->args.hipLaunchKernel.function_address,
data->args.hipLaunchKernel.stream);
tim::auto_lock_t _lk{ tim::type_mutex<data_type_mutex_t>() };
get_roctracer_kernels().emplace(data->correlation_id);
if(get_use_perfetto())
{
get_roctracer_key_map().emplace(data->correlation_id, _name);
}
if(get_use_timemory())
{
get_roctracer_data_map().emplace(
data->correlation_id,
roctracer_bundle_t{ tim::static_string{ _name }() });
}
break;
}
case HIP_API_ID_hipModuleLaunchKernel: {
const char* _name = hipKernelNameRef(data->args.hipModuleLaunchKernel.f);
tim::auto_lock_t _lk{ tim::type_mutex<data_type_mutex_t>() };
get_roctracer_kernels().emplace(data->correlation_id);
if(get_use_perfetto())
{
get_roctracer_key_map().emplace(data->correlation_id, _name);
}
if(get_use_timemory())
{
get_roctracer_data_map().emplace(
data->correlation_id,
roctracer_bundle_t{ tim::static_string{ _name }(),
tim::scope::get_default() });
}
break;
}
default: {
tim::auto_lock_t _lk{ tim::type_mutex<data_type_mutex_t>() };
const char* _name = roctracer_op_string(ACTIVITY_DOMAIN_HIP_API, cid, 0);
if(get_use_perfetto())
{
get_roctracer_key_map().emplace(data->correlation_id, _name);
}
if(get_use_timemory())
{
get_roctracer_data_map().emplace(
data->correlation_id,
roctracer_bundle_t{ tim::static_string{ _name }(),
tim::scope::get_default() });
}
break;
}
}
}
tim::consume_parameters(domain, arg);
}
// Activity tracing callback
void
activity_callback(const char* begin, const char* end, void*)
{
if(!trait::runtime_enabled<comp::roctracer>::get())
return;
static auto _kernel_names = std::unordered_map<const char*, std::string>{};
const roctracer_record_t* record = reinterpret_cast<const roctracer_record_t*>(begin);
const roctracer_record_t* end_record =
reinterpret_cast<const roctracer_record_t*>(end);
std::unordered_set<uint64_t> _indexes{};
tim::auto_lock_t _lk{ tim::type_mutex<data_type_mutex_t>() };
auto& _data = get_roctracer_data_map();
auto& _keys = get_roctracer_key_map();
auto& _kernels = get_roctracer_kernels();
HOSTTRACE_DEBUG("Activity records:\n");
while(record < end_record)
{
HOSTTRACE_DEBUG("\t%-30s\tcorrelation_id(%lu) time_ns(%lu:%lu) device_id(%d) "
"stream_id(%lu)\n",
roctracer_op_string(record->domain, record->correlation_id, 0),
record->correlation_id, record->begin_ns, record->end_ns,
record->device_id, record->queue_id);
auto _is_kernel = _kernels.find(record->correlation_id) != _kernels.end();
if(_is_kernel && record->device_id != 0 && record->queue_id != 0)
{
// these are overheads associated with the kernel launch, not kernel runtime
ROCTRACER_CALL(roctracer_next_record(record, &record));
continue;
}
auto kitr =
(get_use_perfetto()) ? _keys.find(record->correlation_id) : _keys.end();
if(kitr != _keys.end())
{
if(_kernel_names.find(kitr->second) == _kernel_names.end())
_kernel_names.emplace(kitr->second, tim::demangle(kitr->second));
TRACE_EVENT_BEGIN(
"device",
perfetto::StaticString{ _kernel_names.at(kitr->second).c_str() },
record->begin_ns);
TRACE_EVENT_END("device", record->end_ns);
_indexes.emplace(kitr->first);
}
auto itr =
(get_use_timemory()) ? _data.find(record->correlation_id) : _data.end();
if(itr != _data.end())
{
itr->second.start()
.store(std::plus<double>{},
static_cast<double>(record->end_ns - record->begin_ns))
.stop();
_indexes.emplace(kitr->first);
}
// code
ROCTRACER_CALL(roctracer_next_record(record, &record));
}
if(get_use_perfetto())
{
for(auto& itr : _indexes)
_keys.erase(itr);
}
if(get_use_timemory())
{
for(auto& itr : _indexes)
_data.erase(itr);
}
HOSTTRACE_DEBUG("[%s] recorded %lu phases\n", __FUNCTION__,
(unsigned long) _indexes.size());
}
bool&
roctracer_is_setup()
{
static bool _v = false;
return _v;
}
} // namespace
#if !defined(HOSTTRACE_ROCTRACER_LIBKFDWRAPPER)
# define HOSTTRACE_ROCTRACER_LIBKFDWRAPPER "libkfdwrapper64.so"
#endif
namespace tim
{
namespace component
{
void
roctracer::preinit()
{
HOSTTRACE_DEBUG("[%s]\n", __FUNCTION__);
roctracer_data::label() = "roctracer";
roctracer_data::description() = "ROCm tracer (activity API)";
}
void
roctracer::setup()
{
if(!get_use_timemory() && !get_use_perfetto())
return;
auto_lock_t _lk{ type_mutex<roctracer>() };
if(roctracer_is_setup())
return;
roctracer_is_setup() = true;
HOSTTRACE_DEBUG("[%s]\n", __FUNCTION__);
auto libpath = tim::get_env<std::string>("HOSTTRACE_ROCTRACER_LIBKFDWRAPPER",
HOSTTRACE_ROCTRACER_LIBKFDWRAPPER, false);
auto libhandle = dlopen(libpath.c_str(), RTLD_NOW | RTLD_GLOBAL);
if(!libhandle)
fprintf(stderr, "%s\n", dlerror());
dlerror(); // Clear any existing error
// Allocating tracing pool
roctracer_properties_t properties{};
properties.buffer_size = 0x1000;
properties.buffer_callback_fun = activity_callback;
ROCTRACER_CALL(roctracer_open_pool(&properties));
ROCTRACER_CALL(roctracer_set_properties(ACTIVITY_DOMAIN_HIP_API, nullptr));
// Enable API callbacks, all domains
ROCTRACER_CALL(roctracer_enable_callback(hip_api_callback, nullptr));
/*ROCTRACER_CALL(roctracer_enable_op_callback(ACTIVITY_DOMAIN_HIP_API,
HIP_API_ID_hipModuleLaunchKernel,
hip_api_callback, nullptr));
ROCTRACER_CALL(roctracer_enable_op_callback(
ACTIVITY_DOMAIN_HIP_API, HIP_API_ID_hipLaunchKernel, hip_api_callback, nullptr));
ROCTRACER_CALL(roctracer_enable_op_callback(
ACTIVITY_DOMAIN_HIP_API, HIP_API_ID_hipMalloc, hip_api_callback, nullptr));
ROCTRACER_CALL(roctracer_enable_op_callback(
ACTIVITY_DOMAIN_HIP_API, HIP_API_ID_hipMemcpy, hip_api_callback, nullptr));
ROCTRACER_CALL(roctracer_enable_op_callback(
ACTIVITY_DOMAIN_HIP_API, HIP_API_ID_hipFree, hip_api_callback, nullptr));*/
// Enable activity tracing, all domains
ROCTRACER_CALL(roctracer_enable_activity());
if(libhandle)
dlclose(libhandle);
}
void
roctracer::tear_down()
{
auto_lock_t _lk{ type_mutex<roctracer>() };
if(!roctracer_is_setup())
return;
roctracer_is_setup() = false;
HOSTTRACE_DEBUG("[%s]\n", __FUNCTION__);
// flush all the activity
ROCTRACER_CALL(roctracer_flush_activity());
// Disable tracing and closing the pool
ROCTRACER_CALL(roctracer_disable_callback());
/*ROCTRACER_CALL(roctracer_disable_op_callback(ACTIVITY_DOMAIN_HIP_API,
HIP_API_ID_hipModuleLaunchKernel));
ROCTRACER_CALL(roctracer_disable_op_callback(ACTIVITY_DOMAIN_HIP_API,
HIP_API_ID_hipLaunchKernel));
ROCTRACER_CALL(
roctracer_disable_op_callback(ACTIVITY_DOMAIN_HIP_API, HIP_API_ID_hipMemcpy));
ROCTRACER_CALL(
roctracer_disable_op_callback(ACTIVITY_DOMAIN_HIP_API, HIP_API_ID_hipMalloc));
ROCTRACER_CALL(
roctracer_disable_op_callback(ACTIVITY_DOMAIN_HIP_API, HIP_API_ID_hipFree));
*/
ROCTRACER_CALL(roctracer_disable_activity());
ROCTRACER_CALL(roctracer_close_pool());
}
void
roctracer::start()
{
if(tracker_type::start() == 0)
setup();
}
void
roctracer::stop()
{
// flush all the activity
ROCTRACER_CALL(roctracer_flush_activity());
if(tracker_type::stop() == 0)
tear_down();
}
} // namespace component
} // namespace tim
TIMEMORY_INSTANTIATE_EXTERN_COMPONENT(roctracer, false, void)
TIMEMORY_INSTANTIATE_EXTERN_COMPONENT(roctracer_data, true, double)