From a061b7947f3b3b7de18be0b5f1c73ce1cd083f71 Mon Sep 17 00:00:00 2001 From: "Jonathan R. Madsen" Date: Mon, 6 Sep 2021 22:23:24 -0500 Subject: [PATCH] Integrated perfetto + roctracer (#5) - 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 --- .gitmodules | 3 + CMakeLists.txt | 8 + README.md | 88 ++++-- cmake/Compilers.cmake | 8 +- cmake/Findroctracer.cmake | 104 +++++++ cmake/MacroUtilities.cmake | 48 ++++ cmake/Packages.cmake | 57 +++- examples/parallel-overhead/CMakeLists.txt | 11 +- examples/transpose/CMakeLists.txt | 105 ++----- examples/transpose/Makefile | 40 --- examples/transpose/Makefile.in | 39 --- examples/transpose/transpose.cpp | 7 +- external/elfio | 1 + external/timemory | 2 +- include/hosttrace.hpp | 6 - include/library.hpp | 73 ++++- include/roctracer.hpp | 59 ++++ scripts/gperftools-cpu-profile.sh | 138 ++++++++++ scripts/hosttrace_launch_compiler | 75 +++++ src/hosttrace-details.cpp | 2 +- src/hosttrace.cpp | 38 ++- src/library.cpp | 158 ++++++++--- src/roctracer.cpp | 317 ++++++++++++++++++++++ 23 files changed, 1119 insertions(+), 268 deletions(-) create mode 100644 cmake/Findroctracer.cmake delete mode 100644 examples/transpose/Makefile delete mode 100644 examples/transpose/Makefile.in create mode 160000 external/elfio create mode 100644 include/roctracer.hpp create mode 100755 scripts/gperftools-cpu-profile.sh create mode 100755 scripts/hosttrace_launch_compiler create mode 100644 src/roctracer.cpp diff --git a/.gitmodules b/.gitmodules index 2a1df91aa5..43e9a65f6e 100644 --- a/.gitmodules +++ b/.gitmodules @@ -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 diff --git a/CMakeLists.txt b/CMakeLists.txt index bbb503dca5..22a3834704 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -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 $ $ $ diff --git a/README.md b/README.md index ba6786f29e..e57c4bfa5d 100755 --- a/README.md +++ b/README.md @@ -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 -- ``` -### 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 ``` diff --git a/cmake/Compilers.cmake b/cmake/Compilers.cmake index bd655efc46..11ab73d12a 100644 --- a/cmake/Compilers.cmake +++ b/cmake/Compilers.cmake @@ -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() diff --git a/cmake/Findroctracer.cmake b/cmake/Findroctracer.cmake new file mode 100644 index 0000000000..c35c5c9cbb --- /dev/null +++ b/cmake/Findroctracer.cmake @@ -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) + +#------------------------------------------------------------------------------# diff --git a/cmake/MacroUtilities.cmake b/cmake/MacroUtilities.cmake index 2b60617c57..a987a9b638 100644 --- a/cmake/MacroUtilities.cmake +++ b/cmake/MacroUtilities.cmake @@ -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) diff --git a/cmake/Packages.cmake b/cmake/Packages.cmake index 683f7745d9..bae6c16c29 100644 --- a/cmake/Packages.cmake +++ b/cmake/Packages.cmake @@ -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") diff --git a/examples/parallel-overhead/CMakeLists.txt b/examples/parallel-overhead/CMakeLists.txt index bb2660e19e..f4b1698f08 100644 --- a/examples/parallel-overhead/CMakeLists.txt +++ b/examples/parallel-overhead/CMakeLists.txt @@ -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() diff --git a/examples/transpose/CMakeLists.txt b/examples/transpose/CMakeLists.txt index 6d68c7ae3b..2e5ddc7c3d 100644 --- a/examples/transpose/CMakeLists.txt +++ b/examples/transpose/CMakeLists.txt @@ -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() diff --git a/examples/transpose/Makefile b/examples/transpose/Makefile deleted file mode 100644 index f5c2277a0e..0000000000 --- a/examples/transpose/Makefile +++ /dev/null @@ -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.* - diff --git a/examples/transpose/Makefile.in b/examples/transpose/Makefile.in deleted file mode 100644 index d3c62eb59e..0000000000 --- a/examples/transpose/Makefile.in +++ /dev/null @@ -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.* diff --git a/examples/transpose/transpose.cpp b/examples/transpose/transpose.cpp index dce4a6fdec..51aaf340e4 100644 --- a/examples/transpose/transpose.cpp +++ b/examples/transpose/transpose.cpp @@ -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>(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); diff --git a/external/elfio b/external/elfio new file mode 160000 index 0000000000..637221d267 --- /dev/null +++ b/external/elfio @@ -0,0 +1 @@ +Subproject commit 637221d26748171ae2a29d4abea9571036b05d66 diff --git a/external/timemory b/external/timemory index aa4a0ed7b2..cc0cfd6101 160000 --- a/external/timemory +++ b/external/timemory @@ -1 +1 @@ -Subproject commit aa4a0ed7b25e78b5ecc433b38c5ecce72953d614 +Subproject commit cc0cfd610169a65d36b6e9058b9166e7e1647e74 diff --git a/include/hosttrace.hpp b/include/hosttrace.hpp index 824a93c127..1dc41b0407 100644 --- a/include/hosttrace.hpp +++ b/include/hosttrace.hpp @@ -396,9 +396,6 @@ struct module_function std::pair _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; diff --git a/include/library.hpp b/include/library.hpp index 9b9ea67d7e..a8ced84ab8 100644 --- a/include/library.hpp +++ b/include/library.hpp @@ -4,12 +4,14 @@ #if !defined(TIMEMORY_USE_PERFETTO) # include # 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 @@ -122,12 +131,16 @@ private: const char* m_prefix = nullptr; }; +using papi_tot_ins = comp::papi_tuple; using fork_gotcha_t = comp::gotcha<4, tim::component_tuple, hosttrace>; using mpi_gotcha_t = comp::gotcha<4, tim::component_tuple, hosttrace>; using hosttrace_bundle_t = tim::lightweight_tuple; + 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; using bundle_t = tim::component_bundle; using bundle_allocator_t = tim::data::ring_buffer_allocator; @@ -190,6 +203,54 @@ get_state(); std::unique_ptr& get_main_bundle(); +bool +get_use_perfetto(); + +bool +get_use_timemory(); + +//--------------------------------------------------------------------------------------// + +template +struct hosttrace_thread_data +{ + static constexpr size_t max_supported_threads = MaxThreads; + using instance_array_t = std::array, max_supported_threads>; + + template + static void construct(Args&&...); + static std::unique_ptr& instance(); + static instance_array_t& instances(); +}; + +template +template +void +hosttrace_thread_data::construct(Args&&... _args) +{ + static thread_local bool _v = [&_args...]() { + instances().at(threading::get_id()) = + std::make_unique(std::forward(_args)...); + return true; + }(); + (void) _v; +} + +template +std::unique_ptr& +hosttrace_thread_data::instance() +{ + return instances().at(threading::get_id()); +} + +template +typename hosttrace_thread_data::instance_array_t& +hosttrace_thread_data::instances() +{ + static auto _v = instance_array_t{}; + return _v; +} + //--------------------------------------------------------------------------------------// // there are currently some strange things that happen with vector so using diff --git a/include/roctracer.hpp b/include/roctracer.hpp new file mode 100644 index 0000000000..69046ccd73 --- /dev/null +++ b/include/roctracer.hpp @@ -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; + +struct roctracer +: base +, private policy::instance_tracker +{ + using value_type = void; + using base_type = base; + using tracker_type = policy::instance_tracker; + + 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) diff --git a/scripts/gperftools-cpu-profile.sh b/scripts/gperftools-cpu-profile.sh new file mode 100755 index 0000000000..06624c4d4e --- /dev/null +++ b/scripts/gperftools-cpu-profile.sh @@ -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 "${_PATH}" +} + +# 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 diff --git a/scripts/hosttrace_launch_compiler b/scripts/hosttrace_launch_compiler new file mode 100755 index 0000000000..7d09296370 --- /dev/null +++ b/scripts/hosttrace_launch_compiler @@ -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 diff --git a/src/hosttrace-details.cpp b/src/hosttrace-details.cpp index 7f24eaa2be..6eeed9c8cb 100644 --- a/src/hosttrace-details.cpp +++ b/src/hosttrace-details.cpp @@ -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; } diff --git a/src/hosttrace.cpp b/src/hosttrace.cpp index 3f536319e6..56ae45b553 100644 --- a/src/hosttrace.cpp +++ b/src/hosttrace.cpp @@ -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> hash_ids = {}; static std::map use_stubs = {}; static std::map beg_stubs = {}; @@ -282,6 +283,15 @@ main(int argc, char** argv) .count(1) .dtype("size_t") .action([](parser_t& p) { batch_size = p.get("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("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); diff --git a/src/library.cpp b/src/library.cpp index 7e1bc26efe..cb8be98291 100644 --- a/src/library.cpp +++ b/src/library.cpp @@ -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("HOSTTRACE_SAMPLE_RATE", 1); + return _v; +} + +std::vector& +get_sample_data() +{ + static thread_local auto _v = std::vector{}; + 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()->get_is_running()); + auto& _main_bundle = get_main_bundle(); + _main_bundle->start(); + assert(_main_bundle->get()->get_is_running()); +#if defined(HOSTTRACE_USE_ROCTRACER) + assert(_main_bundle->get() != nullptr); + assert(_main_bundle->get()->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::construct( + TIMEMORY_JOIN("", _exe, "/thread-", threading::get_id()), + quirk::config{}); + static thread_local auto _dtor = scope::destructor{ []() { + hosttrace_thread_data::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::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::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(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::instances()) + { + if(itr && itr->get() && + !itr->get()->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(trace_data.size()) / units::KB, + static_cast(trace_data.size()) / units::MB, + static_cast(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; diff --git a/src/roctracer.cpp b/src/roctracer.cpp new file mode 100644 index 0000000000..e8c0b1fe30 --- /dev/null +++ b/src/roctracer.cpp @@ -0,0 +1,317 @@ + +#include "roctracer.hpp" +#include "library.hpp" + +#include +#include +#include + +// 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; + +namespace units = tim::units; + +namespace +{ +static auto& +get_roctracer_kernels() +{ + static auto _v = std::unordered_set{}; + return _v; +} + +static auto& +get_roctracer_data_map() +{ + static auto _v = std::unordered_map{}; + return _v; +} + +static auto& +get_roctracer_key_map() +{ + static auto _v = std::unordered_map{}; + return _v; +} + +using data_type_mutex_t = std::decay_t; + +// HIP API callback function +void +hip_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg) +{ + if(!trait::runtime_enabled::get()) + return; + const hip_api_data_t* data = reinterpret_cast(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() }; + 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() }; + 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() }; + 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::get()) + return; + static auto _kernel_names = std::unordered_map{}; + const roctracer_record_t* record = reinterpret_cast(begin); + const roctracer_record_t* end_record = + reinterpret_cast(end); + std::unordered_set _indexes{}; + + tim::auto_lock_t _lk{ tim::type_mutex() }; + 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{}, + static_cast(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() }; + if(roctracer_is_setup()) + return; + roctracer_is_setup() = true; + HOSTTRACE_DEBUG("[%s]\n", __FUNCTION__); + + auto libpath = tim::get_env("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() }; + 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)