diff --git a/projects/rocprofiler-sdk/.github/workflows/docs.yml b/projects/rocprofiler-sdk/.github/workflows/docs.yml index 3a9460665c..82871c13ed 100644 --- a/projects/rocprofiler-sdk/.github/workflows/docs.yml +++ b/projects/rocprofiler-sdk/.github/workflows/docs.yml @@ -9,7 +9,7 @@ on: - 'VERSION' - 'source/docs/**' - 'source/scripts/update-docs.sh' - - 'source/include/rocprofiler/*' + - 'source/include/rocprofiler-sdk/**' - '.github/workflows/docs.yml' pull_request: branches: [main] @@ -18,7 +18,7 @@ on: - 'VERSION' - 'source/docs/**' - 'source/scripts/update-docs.sh' - - 'source/include/rocprofiler/*' + - 'source/include/rocprofiler-sdk/**' - '.github/workflows/docs.yml' concurrency: diff --git a/projects/rocprofiler-sdk/cmake/Templates/rocprofiler-sdk/config.cmake.in b/projects/rocprofiler-sdk/cmake/Templates/rocprofiler-sdk/config.cmake.in index d89d799d92..eba09ad26c 100644 --- a/projects/rocprofiler-sdk/cmake/Templates/rocprofiler-sdk/config.cmake.in +++ b/projects/rocprofiler-sdk/cmake/Templates/rocprofiler-sdk/config.cmake.in @@ -96,8 +96,7 @@ else() @hip_DIR@ @rocm_version_DIR@) - include( - "${@PACKAGE_NAME@_CMAKE_DIR}/@PACKAGE_NAME@-config-nolink-target.cmake") + include("${@PACKAGE_NAME@_CMAKE_DIR}/@PACKAGE_NAME@-config-nolink-target.cmake") add_library(@PACKAGE_NAME@::@PACKAGE_NAME@-external-nolink INTERFACE IMPORTED) @@ -129,11 +128,15 @@ else() endforeach() endif() - target_link_libraries( - @PROJECT_NAME@::@PROJECT_NAME@ - INTERFACE @PACKAGE_NAME@::@PACKAGE_NAME@-external-nolink) + target_link_libraries(@PROJECT_NAME@::@PROJECT_NAME@ + INTERFACE @PACKAGE_NAME@::@PACKAGE_NAME@-external-nolink) endif() +add_executable(@PACKAGE_NAME@::rocprofv3 IMPORTED) +set_property( + TARGET @PACKAGE_NAME@::rocprofv3 + PROPERTY IMPORTED_LOCATION ${@PACKAGE_NAME@_ROOT_DIR}/@CMAKE_INSTALL_BINDIR@/rocprofv3) + include(FindPackageHandleStandardArgs) find_package_handle_standard_args( @PACKAGE_NAME@ diff --git a/projects/rocprofiler-sdk/cmake/rocprofiler_memcheck.cmake b/projects/rocprofiler-sdk/cmake/rocprofiler_memcheck.cmake index cbd5a52f4b..92a34d6583 100644 --- a/projects/rocprofiler-sdk/cmake/rocprofiler_memcheck.cmake +++ b/projects/rocprofiler-sdk/cmake/rocprofiler_memcheck.cmake @@ -25,7 +25,7 @@ endfunction() function(rocprofiler_set_memcheck_env _TYPE _LIB_BASE) set(_LIBS ${_LIB_BASE}) - foreach(_N 6 5 4 3 2 1 0) + foreach(_N ${ARGN} 6 5 4 3 2 1 0) list( APPEND _LIBS ${CMAKE_SHARED_LIBRARY_PREFIX}${_LIB_BASE}${CMAKE_SHARED_LIBRARY_SUFFIX}.${_N} @@ -34,7 +34,7 @@ function(rocprofiler_set_memcheck_env _TYPE _LIB_BASE) foreach(_LIB ${_LIBS}) if(NOT ${_TYPE}_LIBRARY) - find_library(${_TYPE}_LIBRARY NAMES ${_LIB} ${ARGN}) + find_library(${_TYPE}_LIBRARY NAMES ${_LIB}) endif() endforeach() @@ -60,7 +60,7 @@ elseif(ROCPROFILER_MEMCHECK STREQUAL "MemorySanitizer") rocprofiler_add_memcheck_flags("memory") elseif(ROCPROFILER_MEMCHECK STREQUAL "ThreadSanitizer") rocprofiler_add_memcheck_flags("thread") - rocprofiler_set_memcheck_env("${ROCPROFILER_MEMCHECK}" "tsan") + rocprofiler_set_memcheck_env("${ROCPROFILER_MEMCHECK}" "tsan" 0) elseif(ROCPROFILER_MEMCHECK STREQUAL "UndefinedBehaviorSanitizer") rocprofiler_add_memcheck_flags("undefined") rocprofiler_set_memcheck_env("${ROCPROFILER_MEMCHECK}" "ubsan") diff --git a/projects/rocprofiler-sdk/source/bin/rocprofv3 b/projects/rocprofiler-sdk/source/bin/rocprofv3 index f07702c256..68d8ba6037 100755 --- a/projects/rocprofiler-sdk/source/bin/rocprofv3 +++ b/projects/rocprofiler-sdk/source/bin/rocprofv3 @@ -1,11 +1,13 @@ #!/bin/bash -e -set -eou pipefail +set -eo pipefail -ROCPROFV2_DIR=$(dirname -- "$(realpath "${BASH_SOURCE[0]}")"); -ROCM_DIR=$( dirname -- "${ROCPROFV2_DIR}"; ) +ROCPROFV3_DIR=$(dirname -- "$(realpath "${BASH_SOURCE[0]}")") +ROCM_DIR=$(dirname -- "${ROCPROFV3_DIR}") -export HSA_TOOLS_LIB=${ROCM_DIR}/lib/librocprofiler-sdk.so.0 +: ${HSA_TOOLS_LIB:="${ROCM_DIR}/lib/librocprofiler-sdk.so"} + +export HSA_TOOLS_LIB # Define color codes GREEN='\033[0;32m' @@ -13,61 +15,130 @@ GREY='\033[0;90m' RESET='\033[0m' usage() { + local EC=${1} + if [ -z "${EC}" ]; then EC=1; fi echo -e "${RESET}ROCProfilerV3 Run Script Usage:" echo -e "${GREEN}-h | --help ${RESET} For showing this message" + echo -e "${GREEN}-i | --input ${RESET} For adding counters file path (every line in the text file represents a counter)" echo -e "${GREEN}--hsa-trace ${RESET} For Collecting HSA API Traces" echo -e "${GREEN}--kernel-trace ${RESET} For Collecting Kernel Dispatch Traces" + echo -e "${GREEN}--memory-copy-trace ${RESET} For Collecting Memory Copy Traces" echo -e "${GREEN}-o | --output-file ${RESET} For the output file name" echo -e "\t#${GREY} usage e.g:(with current dir): rocprofv3 --hsa-trace -o " echo -e "\t#${GREY} usage e.g:(with custom dir): rocprofv3 --hsa-trace -d -o ${RESET}\n" echo -e "${GREEN}-d | --output-directory ${RESET} For adding output path where the output files will be saved" - echo -e "\t#${GREY} usage e.g:(with custom dir): rocprofv3 --hsa-trace -d ${RESET}\n" - exit 1 + echo -e "\t#${GREY} usage e.g:(with custom dir): rocprofv3 --hsa-trace -d ${RESET}" + echo -e "${GREEN}-M | --mangled-kernels ${RESET} Do not demangle the kernel names" + echo -e "${GREEN}-T | --truncate-kernels ${RESET} Truncate the demangled kernel names" + echo -e "" + exit ${EC} } if [ -z "$1" ]; then - usage + usage 1 fi -: "${ROCPROFILER_OUTPUT_PATH:="."}" -: "${ROCPROFILER_OUTPUT_FILE_NAME:=""}" +: ${ROCPROF_OUTPUT_PATH:="."} +: ${ROCPROF_OUTPUT_PATH_INTERNAL:="."} +: ${ROCPROF_OUTPUT_FILE_NAME:=""} +: ${ROCPROF_COUNTERS_PATH:=""} +: ${ROCPROF_PRELOAD:=""} while true; do if [[ "$1" == "-h" || "$1" == "--help" ]]; then - usage + usage 0 + elif [[ "$1" == "-M" || "$1" == "--mangled-kernels" ]]; then + export ROCPROF_DEMANGLE_KERNELS=0 + shift + elif [[ "$1" == "-T" || "$1" == "--truncate-kernels" ]]; then + export ROCPROF_TRUNCATE_KERNELS=1 + shift + elif [[ "$1" == "-i" || "$1" == "--input" ]]; then + if [ "$2" ] && [ -n "$2" ] && [ -r "$2" ]; then + export ROCPROF_COUNTERS_PATH=$2 + export ROCPROF_COUNTER_COLLECTION=1 + else + echo -e "Error: input file \"$2\" doesn't exist!" + usage 1 + fi + shift + shift elif [[ "$1" == "-o" || "$1" == "--output-file-name" ]]; then if [ "$2" ]; then - export ROCPROFILER_OUTPUT_FILE_NAME=$2 + export ROCPROF_OUTPUT_FILE_NAME=$2 else - usage + usage 1 fi shift shift elif [[ "$1" == "-d" || "$1" == "--output-directory" ]]; then if [ "$2" ]; then - ROCPROFILER_OUTPUT_PATH=$2 + ROCPROF_OUTPUT_PATH_INTERNAL=$2 + export ROCPROF_OUTPUT_PATH=$ROCPROF_OUTPUT_PATH_INTERNAL else - usage + usage 1 fi shift shift elif [ "$1" == "--hsa-trace" ]; then - export ROCPROFILER_HSA_API_TRACE=1 + export ROCPROF_HSA_API_TRACE=1 shift elif [ "$1" == "--kernel-trace" ]; then - export ROCPROFILER_KERNEL_TRACE=1 + export ROCPROF_KERNEL_TRACE=1 + shift + elif [ "$1" == "--memory-copy-trace" ]; then + export ROCPROF_MEMORY_COPY_TRACE=1 shift elif [ "$1" == "--" ]; then shift break elif [[ "$1" == "-"* || "$1" == "--"* ]]; then echo -e "Wrong option \"$1\", Please use the following options:\n" - usage + usage 1 else break fi done -export ROCPROFILER_OUTPUT_PATH +# read input counter file +PMC_LINES=() +if [ -n "$ROCPROF_COUNTERS_PATH" ]; then + input=$ROCPROF_COUNTERS_PATH + while IFS= read -r line || [[ -n "$line" ]]; do + #skip empty lines + if [[ -z "$line" ]]; then + continue + fi + PMC_LINES+=("$line") + done <"$input" +fi -ROCP_TOOL_LIBRARIES=${ROCM_DIR}/lib/rocprofiler-sdk/librocprofiler-sdk-tool.so "${@}" +if [ -n "${PMC_LINES:-}" ]; then + #for counter collection + COUNTER=1 + for i in "${!PMC_LINES[@]}"; do + export ROCPROF_COUNTERS="${PMC_LINES[$i]}" + if [[ ! ${PMC_LINES[$i]} =~ "pmc" ]]; then + continue + fi + + RESULT_PATH="$ROCPROF_OUTPUT_PATH_INTERNAL/pmc_$COUNTER" + if [ -n "$ROCPROF_OUTPUT_FILE_NAME" ] || [ -n "$ROCPROF_OUTPUT_PATH" ]; then + export ROCPROF_OUTPUT_PATH=$RESULT_PATH + fi + ((COUNTER++)) + ROCP_TOOL_LIBRARIES="${ROCM_DIR}/lib/rocprofiler-sdk/librocprofiler-sdk-tool.so" \ + LD_LIBRARY_PATH=${ROCM_DIR}/lib:${LD_LIBRARY_PATH} \ + LD_PRELOAD="${ROCPROF_PRELOAD}:${ROCM_DIR}/lib/librocprofiler-sdk.so" \ + "${@}" + if [ -n "$ROCPROF_OUTPUT_PATH" ]; then + echo -e "\nThe output path for the following counters: $ROCPROF_OUTPUT_PATH" + fi + done +else + # for non counter collection. e.g: tracing + ROCP_TOOL_LIBRARIES="${ROCM_DIR}/lib/rocprofiler-sdk/librocprofiler-sdk-tool.so" \ + LD_LIBRARY_PATH=${ROCM_DIR}/lib:${LD_LIBRARY_PATH} \ + LD_PRELOAD="${ROCPROF_PRELOAD}:${ROCM_DIR}/lib/librocprofiler-sdk.so" \ + "${@}" +fi diff --git a/projects/rocprofiler-sdk/source/bin/tests/CMakeLists.txt b/projects/rocprofiler-sdk/source/bin/tests/CMakeLists.txt deleted file mode 100644 index 8b13789179..0000000000 --- a/projects/rocprofiler-sdk/source/bin/tests/CMakeLists.txt +++ /dev/null @@ -1 +0,0 @@ - diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h index 63f0329064..ad36cc4a9d 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h @@ -171,7 +171,7 @@ typedef enum // NOLINT(performance-enum-size) */ typedef enum // NOLINT(performance-enum-size) { - ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_NONE = -1, ///< Unknown memory copy direction + ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_NONE = 0, ///< Unknown memory copy direction ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_HOST_TO_HOST, ///< Memory copy from host to host ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_HOST_TO_DEVICE, ///< Memory copy from host to device ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_DEVICE_TO_HOST, ///< Memory copy from device to host diff --git a/projects/rocprofiler-sdk/source/lib/common/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/common/CMakeLists.txt index f57e1e3aa0..49f0c73fe8 100644 --- a/projects/rocprofiler-sdk/source/lib/common/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/common/CMakeLists.txt @@ -3,10 +3,9 @@ # rocprofiler_activate_clang_tidy() -set(common_sources config.cpp environment.cpp demangle.cpp logging.cpp static_object.cpp - utility.cpp xml.cpp) +set(common_sources environment.cpp demangle.cpp logging.cpp static_object.cpp utility.cpp + xml.cpp) set(common_headers - config.hpp defines.hpp environment.hpp demangle.hpp diff --git a/projects/rocprofiler-sdk/source/lib/common/mpl.hpp b/projects/rocprofiler-sdk/source/lib/common/mpl.hpp index f37180028d..50dc8b879c 100644 --- a/projects/rocprofiler-sdk/source/lib/common/mpl.hpp +++ b/projects/rocprofiler-sdk/source/lib/common/mpl.hpp @@ -23,6 +23,8 @@ #pragma once #include +#include +#include #include #include @@ -107,6 +109,18 @@ template struct is_pair : is_pair_impl>>> {}; +template +struct is_string_type_impl +{ + static constexpr auto value = + is_one_of>::value; +}; + +template +struct is_string_type +: is_string_type_impl>>> +{}; + template constexpr bool is_type_complete_v = false; // NOLINT(misc-definitions-in-headers) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt index d1452e381b..a239f13014 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt @@ -3,8 +3,9 @@ # add_library(rocprofiler-sdk-tool SHARED) -target_sources(rocprofiler-sdk-tool PRIVATE helper.hpp helper.cpp tool.cpp - trace_buffer.hpp) +target_sources( + rocprofiler-sdk-tool PRIVATE config.hpp config.cpp csv.hpp helper.hpp helper.cpp + output_file.hpp output_file.cpp tool.cpp) add_subdirectory(plugins) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/README.md b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/README.md index 56b3c86530..258cb03f1b 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/README.md +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/README.md @@ -17,14 +17,14 @@ hsa-api trace. The tool uses the following environment variables to read the user choices. -- `ROCPROFILER_KERNEL_TRACE=1` to obtain kernel trace -- `ROCPROFILER_HSA_API_TRACE=1` to obtain hsa api trace +- `ROCPROF_KERNEL_TRACE=1` to obtain kernel trace +- `ROCPROF_HSA_API_TRACE=1` to obtain hsa api trace The user can also specify the output filename and output file path to which the traces are written to. -- `ROCPROFILER_OUTPUT_PATH=` to set the output directory path -- `ROCPROFILER_OUTPUT_FILE_NAME=` to set the output file name +- `ROCPROF_OUTPUT_PATH=` to set the output directory path +- `ROCPROF_OUTPUT_FILE_NAME=` to set the output file name ## CHANGELOG diff --git a/projects/rocprofiler-sdk/source/lib/common/config.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.cpp similarity index 76% rename from projects/rocprofiler-sdk/source/lib/common/config.cpp rename to projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.cpp index 1665e2ce87..8463755cd0 100644 --- a/projects/rocprofiler-sdk/source/lib/common/config.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.cpp @@ -21,7 +21,8 @@ // THE SOFTWARE. // -#include "lib/common/config.hpp" +#include "config.hpp" + #include "lib/common/defines.hpp" #include "lib/common/demangle.hpp" #include "lib/common/environment.hpp" @@ -36,23 +37,24 @@ #include #include #include +#include #include #include #include namespace rocprofiler { -namespace common +namespace tool { namespace { -std::time_t* launch_time = new std::time_t{std::time(nullptr)}; +auto launch_time = std::make_unique(std::time_t{std::time(nullptr)}); std::string get_local_datetime(const char* dt_format, std::time_t* dt_curr) { char mbstr[512]; - if(!dt_curr) dt_curr = launch_time; + if(!dt_curr) dt_curr = launch_time.get(); if(std::strftime(mbstr, sizeof(mbstr), dt_format, std::localtime(dt_curr)) != 0) return std::string{mbstr}; @@ -109,6 +111,95 @@ get_num_siblings(pid_t _id = getppid()) { return get_siblings(_id).size(); } + +// replace unsuported specail chars with space +void +handle_special_chars(std::string& str) +{ + // Iterate over the string and replace any special characters with a space. + auto pos = std::string::npos; + while((pos = str.find_first_of("!@#$%&(),*+-./;<=>?@{}^`~|:")) != std::string::npos) + str.at(pos) = ' '; +} + +bool +has_kernel_name_format(std::string const& str) +{ + return std::find_if(str.begin(), str.end(), [](unsigned char ch) { + return (isalnum(ch) != 0 || ch == '_'); + }) != str.end(); +} + +bool +has_counter_format(std::string const& str) +{ + return std::find_if(str.begin(), str.end(), [](unsigned char ch) { + return (isalnum(ch) != 0 || ch == '_'); + }) != str.end(); +} + +// validate kernel names +auto +parse_kernel_names(const std::string& line) +{ + auto kernel_names_v = std::vector{}; + if(line.empty()) return kernel_names_v; + + auto kernel_names = std::set{}; + trim(line); + auto input_line = std::stringstream{line}; + auto kernel_name = std::string{}; + while(getline(input_line, kernel_name, ',')) + { + if(has_kernel_name_format(kernel_name)) + { + LOG(INFO) << "kernel name " << kernel_names.size() << ": " << kernel_name; + kernel_names.emplace(kernel_name); + } + else + { + LOG(ERROR) << "invalid kernel name: " << kernel_name; + } + } + + kernel_names_v.reserve(kernel_names.size()); + for(const auto& itr : kernel_names) + kernel_names_v.emplace_back(itr); + + return kernel_names_v; +} + +std::set +parse_counters(std::string line) +{ + auto counters = std::set{}; + + if(line.empty()) return counters; + + // trim line for any white spaces + trim(line); + + if(!(line[0] == '#' || line.find("pmc") == std::string::npos)) + { + handle_special_chars(line); + + std::stringstream input_line(line); + std::string counter; + while(getline(input_line, counter, ' ')) + { + if(counter.substr(0, 3) != "pmc" && has_counter_format(counter)) + { + counters.emplace(counter); + } + else + { + LOG(ERROR) << "invalid counter: " << counter; + } + } + } + + return counters; +} } // namespace int @@ -127,12 +218,17 @@ get_mpi_rank() return _v; } +config::config() +: kernel_names{parse_kernel_names(get_env("ROCPROF_KERNEL_NAMES", std::string{}))} +, counters{parse_counters(get_env("ROCPROF_COUNTERS", std::string{}))} +{} + std::vector output_keys(std::string _tag) { using strpair_t = std::pair; - auto _cmdline = read_command_line(getpid()); + auto _cmdline = common::read_command_line(getpid()); if(_tag.empty() && !_cmdline.empty()) _tag = ::basename(_cmdline.front().c_str()); @@ -175,11 +271,11 @@ output_keys(std::string _tag) } } - auto* _launch_time = launch_time; + auto* _launch_time = launch_time.get(); auto _time_format = get_env("ROCP_TIME_FORMAT", "%F_%H.%M"); - auto _mpi_size = get_env("OMPI_COMM_WORLD_SIZE", get_env("MV2_COMM_WORLD_SIZE", 0)); - auto _mpi_rank = get_env("OMPI_COMM_WORLD_RANK", get_env("MV2_COMM_WORLD_RANK", -1)); + auto _mpi_size = get_mpi_size(); + auto _mpi_rank = get_mpi_rank(); auto _dmp_size = fmt::format("{}", (_mpi_size) > 0 ? _mpi_size : 1); auto _dmp_rank = fmt::format("{}", (_mpi_rank) > 0 ? _mpi_rank : 0); @@ -312,69 +408,18 @@ format(std::string _fpath, const std::string& _tag) return _fpath; } -std::string -compose_filename(const config& _cfg) -{ - auto _output_path = _cfg.output_path; - auto _output_file = _cfg.output_file; - auto _output_ext = _cfg.output_ext; - - if(_output_path.empty()) _output_path = "."; - if(_cfg.mpi_size > 0) - { - if(_cfg.mpi_rank >= 0) - { - _output_file = fmt::format("{}.{}", _output_file, _cfg.mpi_rank); - } - else - { - _output_file = fmt::format("{}.{}", _output_file, getpid()); - } - } - if(!_output_ext.empty()) - { - if(_output_ext.find('.') == std::string::npos) _output_ext.insert(0, "."); - if(_output_file.length() < _output_ext.length() || - _output_file.find(_output_ext) != _output_file.length() - _output_ext.length()) - _output_file += _output_ext; - } - - // join / and replace any keys with values - auto _prefix = format(common::filesystem::path{_output_path} / _output_file); - - // return on empty - if(_prefix.empty()) return std::string{}; - - // get the absolute path - auto _fname = common::filesystem::absolute(common::filesystem::path{_prefix}); - - // create the directory if necessary - auto _fname_path = _fname.parent_path(); - if(!common::filesystem::exists(_fname_path)) - common::filesystem::create_directories(_fname.parent_path()); - - return _fname.string(); -} - std::string format_name(std::string_view _name, const config& _cfg) { - if(_cfg.demangle && _cfg.truncate) - { - return truncate_name(cxx_demangle(_name)); - } + if(!_cfg.demangle && !_cfg.truncate) return std::string{_name}; - if(_cfg.demangle) - { - return cxx_demangle(_name); - } + // truncating requires demangling first so always demangle + auto _demangled_name = + common::cxx_demangle(std::regex_replace(_name.data(), std::regex{"(\\.kd)$"}, "")); - if(_cfg.truncate) - { - return truncate_name(_name); - } + if(_cfg.truncate) return common::truncate_name(_demangled_name); - return std::string{_name}; + return _demangled_name; } void @@ -388,5 +433,5 @@ output_key::output_key(std::string _key, std::string _val, std::string _desc) , value{std::move(_val)} , description{std::move(_desc)} {} -} // namespace common +} // namespace tool } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/common/config.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp similarity index 67% rename from projects/rocprofiler-sdk/source/lib/common/config.hpp rename to projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp index 84ea6672ef..ad007539a6 100644 --- a/projects/rocprofiler-sdk/source/lib/common/config.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp @@ -24,14 +24,19 @@ #pragma once #include "lib/common/environment.hpp" +#include "lib/common/filesystem.hpp" +#include #include #include namespace rocprofiler { -namespace common +namespace tool { +namespace fs = common::filesystem; +using common::get_env; + enum class config_context { global = 0, @@ -50,13 +55,21 @@ get_mpi_rank(); struct config { - bool demangle = get_env("ROCP_DEMANGLE_KERNELS", true); - bool truncate = get_env("ROCP_TRUNCATE_KERNELS", false); - int mpi_size = get_mpi_size(); - int mpi_rank = get_mpi_rank(); - std::string output_path = get_env("ROCP_OUTPUT_PATH", "."); - std::string output_file = get_env("ROCP_OUTPUT_FILE", "results"); - std::string output_ext = {}; + config(); + + bool demangle = get_env("ROCPROF_DEMANGLE_KERNELS", true); + bool truncate = get_env("ROCPROF_TRUNCATE_KERNELS", false); + bool kernel_trace = get_env("ROCPROF_KERNEL_TRACE", false); + bool hsa_api_trace = get_env("ROCPROF_HSA_API_TRACE", false); + bool marker_api_trace = get_env("ROCPROF_MARKER_API_TRACE", false); + bool memory_copy_trace = get_env("ROCPROF_MEMORY_COPY_TRACE", false); + bool counter_collection = get_env("ROCPROF_COUNTER_COLLECTION", false); + int mpi_size = get_mpi_size(); + int mpi_rank = get_mpi_rank(); + std::string output_path = get_env("ROCPROF_OUTPUT_PATH", fs::current_path().string()); + std::string output_file = get_env("ROCPROF_OUTPUT_FILE_NAME", std::to_string(getpid())); + std::vector kernel_names = {}; + std::set counters = {}; }; template @@ -90,9 +103,6 @@ struct output_key std::vector output_keys(std::string _tag = {}); -std::string -compose_filename(const config&); - std::string format(std::string _fpath, const std::string& _tag = {}); @@ -101,5 +111,5 @@ format_name(std::string_view _name, const config& = get_config<>()); void initialize(); -} // namespace common +} // namespace tool } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/csv.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/csv.hpp new file mode 100644 index 0000000000..c2a79da342 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/csv.hpp @@ -0,0 +1,83 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#include "lib/common/mpl.hpp" + +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace tool +{ +namespace csv +{ +template +std::ostream& +write_csv_entry(std::ostream& ofs, TupleT&& _data, std::index_sequence) +{ + auto _write = [&ofs](size_t idx, auto&& _val) { + if(idx > 0) ofs << ","; + if constexpr(rocprofiler::common::mpl::is_string_type::value) ofs << "\""; + ofs << _val; + if constexpr(rocprofiler::common::mpl::is_string_type::value) ofs << "\""; + }; + + (_write(Idx, std::get(_data)), ...); + return (ofs << '\n'); +} + +template +struct csv_encoder +{ + static constexpr auto columns = NumCols; + + template = 0> + static auto write_row(std::ostream& ofs, Args&&... args) + { + write_csv_entry( + ofs, std::make_tuple(std::forward(args)...), std::make_index_sequence{}); + return csv_encoder{}; + } + + template + static auto write_row(std::ostream& ofs, const std::array& arr) + { + static_assert(N == columns, "Error! too many/few args passed"); + write_csv_entry(ofs, arr, std::make_index_sequence{}); + return csv_encoder{}; + } +}; + +using hsa_csv_encoder = csv_encoder<7>; +using kernel_trace_csv_encoder = csv_encoder<16>; +using counter_collection_csv_encoder = csv_encoder<14>; +using memory_copy_csv_encoder = csv_encoder<7>; +} // namespace csv +} // namespace tool +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.cpp index 0749c9a125..1c7c5c9089 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.cpp @@ -21,6 +21,8 @@ // SOFTWARE. #include "helper.hpp" +#include "config.hpp" +#include "rocprofiler-sdk/fwd.h" #include @@ -28,6 +30,8 @@ #include #include #include +#include +#include namespace { @@ -73,7 +77,6 @@ enum amd_kernel_code_property_t AMD_HSA_BITS_CREATE_ENUM_ENTRIES(AMD_KERNEL_CODE_PROPERTY_RESERVED1, 12, 4), }; -std::mutex kernel_descriptor_name_map_mutex; std::unordered_map kernel_descriptor_name_map; std::mutex kernel_properties_correlation_mutex; @@ -120,6 +123,7 @@ sgpr_count(const std::string_view& name, const kernel_descriptor_t& kernel_code) // GFX10 and later always allocate 128 sgprs. // TODO(srnagara): Recheck the extraction of gfxip from gpu name + const char* name_data = name.data(); const size_t gfxip_label_len = std::min(name.size() - 2, size_t{63}); if(gfxip_label_len > 0 && strnlen(name_data, gfxip_label_len + 1) >= gfxip_label_len) @@ -169,19 +173,12 @@ GetKernelCode(uint64_t kernel_object) } } // namespace -void -SetKernelDescriptorName(rocprofiler_address_t kernel_descriptor, const char* kernel_name) -{ - std::lock_guard kernel_descriptor_name_map_lock(kernel_descriptor_name_map_mutex); - kernel_descriptor_name_map[kernel_descriptor] = kernel_name; -} - void SetKernelProperties(uint64_t correlation_id, rocprofiler_tool_kernel_properties_t kernel_properties) { std::lock_guard kernel_properties_correlation_map_lock( kernel_properties_correlation_mutex); - kernel_properties_correlation_map[correlation_id] = kernel_properties; + kernel_properties_correlation_map[correlation_id] = std::move(kernel_properties); } rocprofiler_tool_kernel_properties_t @@ -197,74 +194,24 @@ GetKernelProperties(uint64_t correlation_id) } return it->second; } -const char* -GetKernelDescriptorName(rocprofiler_address_t kernel_descriptor) -{ - std::lock_guard kernel_descriptor_name_map_lock(kernel_descriptor_name_map_mutex); - auto it = kernel_descriptor_name_map.find(kernel_descriptor); - if(it == kernel_descriptor_name_map.end()) - { - std::cout << "kernel name not found" << std::endl; - abort(); - } - return it->second; -} - -std::vector -GetCounterNames() -{ - std::vector counters; - const char* line_c_str = getenv("ROCPROFILER_COUNTERS"); - if(line_c_str) - { - std::string line = line_c_str; - // skip commented lines - auto found = line.find_first_not_of(" \t"); - if(found != std::string::npos) - { - if(line[found] == '#') return {}; - } - if(line.find("pmc") == std::string::npos) return counters; - char seperator = ' '; - std::string::size_type prev_pos = 0, pos = line.find(seperator, prev_pos); - prev_pos = ++pos; - if(pos != std::string::npos) - { - while((pos = line.find(seperator, pos)) != std::string::npos) - { - std::string substring(line.substr(prev_pos, pos - prev_pos)); - if(substring.length() > 0 && substring != ":") - { - counters.push_back(substring); - } - prev_pos = ++pos; - } - if(!line.substr(prev_pos, pos - prev_pos).empty()) - { - counters.push_back(line.substr(prev_pos, pos - prev_pos)); - } - } - } - return counters; -} void populate_kernel_properties_data(rocprofiler_tool_kernel_properties_t* kernel_properties, - const hsa_kernel_dispatch_packet_t dispatch_packet) + const hsa_kernel_dispatch_packet_t* dispatch_packet) { - const uint64_t kernel_object = dispatch_packet.kernel_object; + const uint64_t kernel_object = dispatch_packet->kernel_object; const kernel_descriptor_t* kernel_code = GetKernelCode(kernel_object); uint64_t grid_size = - dispatch_packet.grid_size_x * dispatch_packet.grid_size_y * dispatch_packet.grid_size_z; + dispatch_packet->grid_size_x * dispatch_packet->grid_size_y * dispatch_packet->grid_size_z; if(grid_size > UINT32_MAX) abort(); kernel_properties->grid_size = grid_size; - uint64_t workgroup_size = dispatch_packet.workgroup_size_x * dispatch_packet.workgroup_size_y * - dispatch_packet.workgroup_size_z; + uint64_t workgroup_size = dispatch_packet->workgroup_size_x * + dispatch_packet->workgroup_size_y * dispatch_packet->workgroup_size_z; if(workgroup_size > UINT32_MAX) abort(); kernel_properties->workgroup_size = (uint32_t) workgroup_size; - kernel_properties->lds_size = dispatch_packet.group_segment_size; - kernel_properties->scratch_size = dispatch_packet.private_segment_size; + kernel_properties->lds_size = dispatch_packet->group_segment_size; + kernel_properties->scratch_size = dispatch_packet->private_segment_size; kernel_properties->arch_vgpr_count = arch_vgpr_count(kernel_properties->gpu_agent.name, *kernel_code); kernel_properties->accum_vgpr_count = @@ -275,79 +222,62 @@ populate_kernel_properties_data(rocprofiler_tool_kernel_properties_t* kernel_pro AMD_KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32) ? 32 : 64; - kernel_properties->signal_handle = dispatch_packet.completion_signal.handle; + kernel_properties->signal_handle = dispatch_packet->completion_signal.handle; } -std::string -cxa_demangle(std::string_view _mangled_name, int* _status) +rocprofiler_tool_buffer_name_info_t +get_buffer_id_names() { - constexpr size_t buffer_len = 4096; - // return the mangled since there is no buffer - if(_mangled_name.empty()) - { - *_status = -2; - return std::string{}; - } + static auto supported = std::unordered_set{ + ROCPROFILER_BUFFER_TRACING_HSA_API, + ROCPROFILER_BUFFER_TRACING_MEMORY_COPY, + ROCPROFILER_BUFFER_TRACING_MARKER_API}; - auto _demangled_name = std::string{_mangled_name}; + auto cb_name_info = rocprofiler_tool_buffer_name_info_t{}; + // + // callback for each kind operation + // + static auto tracing_kind_operation_cb = + [](rocprofiler_buffer_tracing_kind_t kindv, uint32_t operation, void* data_v) { + auto* name_info_v = static_cast(data_v); - // PARAMETERS to __cxa_demangle - // mangled_name: - // A nullptr-terminated character string containing the name to be demangled. - // buffer: - // A region of memory, allocated with malloc, of *length bytes, into which the - // demangled name is stored. If output_buffer is not long enough, it is expanded - // using realloc. output_buffer may instead be nullptr; in that case, the demangled - // name is placed in a region of memory allocated with malloc. - // _buflen: - // If length is non-nullptr, the length of the buffer containing the demangled name - // is placed in *length. - // status: - // *status is set to one of the following values - size_t _demang_len = 0; - char* _demang = abi::__cxa_demangle(_demangled_name.c_str(), nullptr, &_demang_len, _status); - switch(*_status) - { - // 0 : The demangling operation succeeded. - // -1 : A memory allocation failure occurred. - // -2 : mangled_name is not a valid name under the C++ ABI mangling rules. - // -3 : One of the arguments is invalid. - case 0: + if(supported.count(kindv) > 0) + { + const char* name = nullptr; + ROCPROFILER_CALL(rocprofiler_query_buffer_tracing_kind_operation_name( + kindv, operation, &name, nullptr), + "query buffer failed"); + if(name) name_info_v->operation_names[kindv][operation] = name; + } + + return 0; + }; + + // + // callback for each kind (i.e. domain) + // + static auto tracing_kind_cb = [](rocprofiler_buffer_tracing_kind_t kind, void* data) { + // store the buffer kind name + auto* name_info_v = static_cast(data); + const char* name = nullptr; + ROCPROFILER_CALL(rocprofiler_query_buffer_tracing_kind_name(kind, &name, nullptr), + "query buffer failed"); + + if(name) name_info_v->kind_names[kind] = name; + + if(supported.count(kind) > 0) { - if(_demang) _demangled_name = std::string{_demang}; - break; + ROCPROFILER_CALL(rocprofiler_iterate_buffer_tracing_kind_operations( + kind, tracing_kind_operation_cb, static_cast(data)), + "query buffer failed"); } - case -1: - { - char _msg[buffer_len]; - ::memset(_msg, '\0', buffer_len * sizeof(char)); - ::snprintf(_msg, - buffer_len, - "memory allocation failure occurred demangling %s", - _demangled_name.c_str()); - ::perror(_msg); - break; - } - case -2: break; - case -3: - { - char _msg[buffer_len]; - ::memset(_msg, '\0', buffer_len * sizeof(char)); - ::snprintf(_msg, - buffer_len, - "Invalid argument in: (\"%s\", nullptr, nullptr, %p)", - _demangled_name.c_str(), - (void*) _status); - ::perror(_msg); - break; - } - default: break; + + return 0; }; - // if it "demangled" but the length is zero, set the status to -2 - if(_demang_len == 0 && *_status == 0) *_status = -2; + ROCPROFILER_CALL(rocprofiler_iterate_buffer_tracing_kinds(tracing_kind_cb, + static_cast(&cb_name_info)), + "iterate_buffer failed"); - // free allocated buffer - ::free(_demang); - return _demangled_name; + return cb_name_info; } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.hpp index 497115bb41..8011123973 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/helper.hpp @@ -20,6 +20,22 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE // SOFTWARE. +#pragma once + +#include "lib/common/defines.hpp" +#include "lib/common/filesystem.hpp" + +#include +#include + +#include +#include +#include +#include +#include +#include +#include + #include #include #include @@ -29,32 +45,22 @@ #include #include #include +#include #include #include #include #include -#include "lib/common/filesystem.hpp" - -#include -#include -#include -#include -#include -#include -#include -#include -#include - #define ROCPROFILER_CALL(result, msg) \ { \ - rocprofiler_status_t CHECKSTATUS = result; \ - if(CHECKSTATUS != ROCPROFILER_STATUS_SUCCESS) \ + rocprofiler_status_t ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__) = result; \ + if(ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__) != ROCPROFILER_STATUS_SUCCESS) \ { \ - std::string status_msg = rocprofiler_get_status_string(CHECKSTATUS); \ + std::string status_msg = \ + rocprofiler_get_status_string(ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__)); \ std::cerr << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg \ - << " failed with error code " << CHECKSTATUS << ": " << status_msg \ - << std::endl; \ + << " failed with error code " << ROCPROFILER_VARIABLE(CHECKSTATUS, __LINE__) \ + << ": " << status_msg << std::endl; \ std::stringstream errmsg{}; \ errmsg << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg " failure (" \ << status_msg << ")"; \ @@ -86,16 +92,12 @@ typedef struct uint64_t signal_handle; uint64_t kernel_object; rocprofiler_queue_id_t queue_id; + std::string kernel_name; rocprofiler_agent_t gpu_agent; + uint64_t thread_id; } rocprofiler_tool_kernel_properties_t; -typedef struct -{ - std::vector gpu_agents_lists; - -} rocprofiler_tool_agent_callback_t; - struct kernel_descriptor_t { uint8_t reserved0[16]; @@ -108,19 +110,20 @@ struct kernel_descriptor_t uint8_t reserved2[6]; }; -using rocprofiler_tool_callback_kind_names_t = - std::map; -using rocprofiler_tool_callback_kind_operation_names_t = - std::map>; +using rocprofiler_tool_buffer_kind_names_t = + std::unordered_map; +using rocprofiler_tool_buffer_kind_operation_names_t = + std::unordered_map>; -struct rocprofiler_tool_callback_name_info_t +struct rocprofiler_tool_buffer_name_info_t { - rocprofiler_tool_callback_kind_names_t kind_names = {}; - rocprofiler_tool_callback_kind_operation_names_t operation_names = {}; + rocprofiler_tool_buffer_kind_names_t kind_names = {}; + rocprofiler_tool_buffer_kind_operation_names_t operation_names = {}; }; -std::vector -GetCounterNames(); +// std::vector +// GetCounterNames(); void SetKernelDescriptorName(rocprofiler_address_t kernel_descriptor, const char* name); @@ -140,10 +143,7 @@ GetKernelDescriptorName(rocprofiler_address_t kernel_descriptor); void populate_kernel_properties_data(rocprofiler_tool_kernel_properties_t* kernel_properties, - const hsa_kernel_dispatch_packet_t dispatch_packet); + const hsa_kernel_dispatch_packet_t* dispatch_packet); -void -TracerFlushRecord(void* data, rocprofiler_callback_tracing_kind_t kind); - -std::string -cxa_demangle(std::string_view _mangled_name, int* _status); +rocprofiler_tool_buffer_name_info_t +get_buffer_id_names(); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.cpp new file mode 100644 index 0000000000..77945d3382 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.cpp @@ -0,0 +1,62 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include "output_file.hpp" +#include "config.hpp" + +#include + +namespace rocprofiler +{ +namespace tool +{ +namespace fs = common::filesystem; + +std::pair +get_output_stream(const std::string& fname, const std::string& ext) +{ + auto output_path = fs::path{tool::format(tool::get_config().output_path)}; + auto output_file_name = tool::format(tool::get_config().output_file); + + if(output_path.string().empty()) return {&std::clog, [](auto*&) {}}; + + if(fs::exists(output_path) && !fs::is_directory(fs::status(output_path))) + throw std::runtime_error{ + fmt::format("ROCPROFILER_OUTPUT_PATH ({}) already exists and is not a directory", + output_path.string())}; + if(!fs::exists(output_path)) fs::create_directories(output_path); + + auto output_file = tool::format(output_path / (output_file_name + "_" + fname + ext)); + auto* _ofs = new std::ofstream{output_file}; + if(!_ofs && !*_ofs) + throw std::runtime_error{fmt::format("Failed to open {} for output", output_file)}; + + LOG(ERROR) << "Results File: " << output_file; + + return {_ofs, [](std::ostream*& v) { + if(v) dynamic_cast(v)->close(); + delete v; + v = nullptr; + }}; +} +} // namespace tool +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.hpp new file mode 100644 index 0000000000..f49c9b9f96 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.hpp @@ -0,0 +1,101 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#include "config.hpp" +#include "csv.hpp" + +#include "lib/common/filesystem.hpp" +#include "lib/rocprofiler-sdk-tool/csv.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace tool +{ +std::pair +get_output_stream(const std::string& fname, const std::string& ext = ".csv"); + +struct output_file +{ + template + output_file(std::string name, csv::csv_encoder, std::array&& header); + + ~output_file() { m_dtor(m_stream); } + + output_file(const output_file&) = delete; + output_file& operator=(const output_file&) = delete; + + std::string name() const { return m_name; } + + template + std::ostream& operator<<(T&& value) + { + auto _lk = std::unique_lock{m_mutex}; + return ((m_stream) ? *m_stream : std::cerr) << std::forward(value); + } + + std::ostream& operator<<(std::ostream& (*func)(std::ostream&) ) + { + auto _lk = std::unique_lock{m_mutex}; + return ((m_stream) ? *m_stream : std::cerr) << func; + } + + operator bool() const { return m_stream != nullptr; } + +private: + using stream_dtor_t = void (*)(std::ostream*&); + + const std::string m_name = {}; + std::mutex m_mutex = {}; + std::ostream* m_stream = nullptr; + stream_dtor_t m_dtor = [](std::ostream*&) {}; +}; + +template +output_file::output_file(std::string name, + csv::csv_encoder encoder, + std::array&& header) +: m_name{std::move(name)} +{ + std::tie(m_stream, m_dtor) = get_output_stream(m_name); + + for(auto& itr : header) + { + LOG_IF(FATAL, itr.empty()) + << "CSV file for " << m_name << " was not provided the correct number of headers"; + } + + // write the csv header + if(m_stream) encoder.write_row(*m_stream, header); +} +} // namespace tool +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp index 79834a55fd..9dfaa7a9bc 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -20,192 +20,186 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE // SOFTWARE. +#include "config.hpp" +#include "csv.hpp" #include "helper.hpp" -#include "rocprofiler-sdk/context.h" -#include "trace_buffer.hpp" +#include "output_file.hpp" +#include "lib/common/demangle.hpp" #include "lib/common/environment.hpp" #include "lib/common/filesystem.hpp" +#include "lib/common/logging.hpp" +#include "lib/common/synchronized.hpp" +#include "lib/common/utility.hpp" + +#include +#include +#include + +#include #include #include - +#include #include #include #include +#include #include #include +#include #include -namespace common = ::rocprofiler::common; -namespace fs = common::filesystem; - -TRACE_BUFFER_INSTANTIATE(); - -// static const uint32_t lds_block_size = 128 * 4; +namespace common = ::rocprofiler::common; +namespace tool = ::rocprofiler::tool; +static const uint32_t lds_block_size = 128 * 4; namespace -{ -auto tool_buffer = rocprofiler_buffer_id_t{}; -auto context_id = rocprofiler_context_id_t{}; -auto output_path = - fs::path{common::get_env("ROCPROFILER_OUTPUT_PATH", fs::current_path().string())}; -auto output_file_name = - common::get_env("ROCPROFILER_OUTPUT_FILE_NAME", std::to_string(getpid())); - -std::pair -get_output_stream(const std::string& fname, const std::string& ext = ".csv") -{ - if(output_path.string().empty()) return {&std::clog, [](auto*&) {}}; - - if(fs::exists(output_path) && !fs::is_directory(fs::status(output_path))) - throw std::runtime_error{ - fmt::format("ROCPROFILER_OUTPUT_PATH ({}) already exists and is not a directory", - output_path.string())}; - if(!fs::exists(output_path)) fs::create_directories(output_path); - - auto output_file = output_path / (output_file_name + "_" + fname + ext); - auto* _ofs = new std::ofstream{output_file}; - if(!_ofs && !*_ofs) - throw std::runtime_error{fmt::format("Failed to open {} for output", output_file.string())}; - std::cout << "Results File: " << output_file << std::endl; - return {_ofs, [](std::ostream*& v) { - if(v) dynamic_cast(v)->close(); - delete v; - v = nullptr; - }}; -} - -template -std::string -as_hex(Tp _v, size_t _width = 16) -{ - auto _ss = std::stringstream{}; - _ss.fill('0'); - _ss << "0x" << std::hex << std::setw(_width) << _v; - return _ss.str(); -} -} // namespace - -struct output_file -{ - output_file(std::string name, std::vector&& header) - : m_name{std::move(name)} - { - std::tie(m_stream, m_dtor) = get_output_stream(m_name); - auto ss = std::stringstream{}; - for(auto&& itr : header) - { - ss << "," << itr; - } - - // write the csv header - if(!ss.str().empty()) *m_stream << ss.str().substr(1) << '\n'; - } - - ~output_file() { m_dtor(m_stream); } - - output_file(const output_file&) = delete; - output_file& operator=(const output_file&) = delete; - - std::string name() const { return m_name; } - - template - std::ostream& operator<<(T&& value) - { - return (*m_stream) << std::forward(value); - } - - std::ostream& operator<<(std::ostream& (*func)(std::ostream&) ) { return (*m_stream) << func; } - - operator bool() const { return m_stream != nullptr; } - -private: - using stream_dtor_t = void (*)(std::ostream*&); - - const std::string m_name = {}; - std::ostream* m_stream = nullptr; - stream_dtor_t m_dtor = [](std::ostream*&) {}; -}; +{} // namespace auto& get_hsa_api_file() { - static auto _v = - output_file{"hsa_api_trace", - {"DOMAIN", "FUNCTION", "START_TIMESTAMP", "END_TIMESTAMP", "CORRELATION_ID"}}; + static auto _v = tool::output_file{"hsa_api_trace", + tool::csv::hsa_csv_encoder{}, + {"Domain", + "Function", + "Process_Id", + "Thread_Id", + "Correlation_Id", + "Start_Timestamp", + "End_Timestamp"}}; return _v; } auto& get_kernel_trace_file() { - static auto _v = output_file{"kernel_trace", - {"AGENT_ID", - "QUEUE_ID", - "KERNEL_ID", - "KERNEL_NAME", - "CONTEXT_ID", - "BUFFER_ID", - "CORRELATION_ID", - "KIND", - "START_TS", - "END_TS", - "PRIVATE_SEGMENT_SIZE", - "GROUP_SEGMENT_SIZE", - "WORKGROUP_SIZE_X", - "WORKGROUP_SIZE_Y", - "WORKGROUP_SIZE_Z", - "GRID_SIZE_X", - "GRID_SIZE_Y", - "GRID_SIZE_Z"}}; + static auto _v = tool::output_file{"kernel_trace", + tool::csv::kernel_trace_csv_encoder{}, + {"Kind", + "Agent_Id", + "Queue_Id", + "Kernel_Id", + "Kernel_Name", + "Correlation_Id", + "Start_Timestamp", + "End_Timestamp", + "Private_Segment_Size", + "Group_Segment_Size", + "Workgroup_Size_X", + "Workgroup_Size_Y", + "Workgroup_Size_Z", + "Grid_Size_X", + "Grid_Size_Y", + "Grid_Size_Z"}}; return _v; } -std::shared_mutex kernel_data_mutex; -std::unordered_map kernel_data; - -struct hsa_api_trace_entry_t +auto& +get_counter_collection_file() { - std::atomic valid; - rocprofiler_callback_tracing_record_t record; - rocprofiler_timestamp_t begin_timestamp; - rocprofiler_timestamp_t end_timestamp; - std::string_view api_name; + static auto _v = tool::output_file{"counter_collection", + tool::csv::counter_collection_csv_encoder{}, + {"Counter_Id", + "Agent_Id", + "Queue_Id", + "Process_Id", + "Thread_Id", + "Grid_Size", + "Kernel-Name", + "Workgroup_Size", + "LDS_Block_Size", + "Scratch_Size", + "VGPR_Count", + "SGPR_Count", + "Counter_Name", + "Counter_Value"}}; + return _v; +} - hsa_api_trace_entry_t(rocprofiler_timestamp_t begin, - rocprofiler_timestamp_t end, - rocprofiler_callback_tracing_record_t tracer_record, - std::string_view name) - : valid(TRACE_ENTRY_INIT) - , record(tracer_record) - , begin_timestamp(begin) - , end_timestamp(end) - , api_name(name) +auto& +get_memory_copy_trace_file() +{ + static auto _v = tool::output_file{"memory_copy_trace", + tool::csv::memory_copy_csv_encoder{}, + {"Kind", + "Direction", + "Source_Agent_Id", + "Destination_Agent_Id", + "Correlation_Id", + "Start_Timestamp", + "End_Timestamp"}}; + return _v; +} + +rocprofiler_buffer_id_t& +get_hsa_api_trace_buffer() +{ + static rocprofiler_buffer_id_t hsa_api_buf = {}; + return hsa_api_buf; +} + +rocprofiler_buffer_id_t& +get_kernel_trace_buffer() +{ + static rocprofiler_buffer_id_t kernel_trace_buf = {}; + return kernel_trace_buf; +} + +rocprofiler_buffer_id_t& +get_counter_collection_buffer() +{ + static rocprofiler_buffer_id_t counter_collection_buf = {}; + return counter_collection_buf; +} + +rocprofiler_buffer_id_t& +get_memory_copy_trace_buffer() +{ + static rocprofiler_buffer_id_t memory_copy_buf = {}; + return memory_copy_buf; +} + +using rocprofiler_kernel_symbol_data_t = + rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t; + +struct kernel_symbol_data : rocprofiler_kernel_symbol_data_t +{ + using base_type = rocprofiler_kernel_symbol_data_t; + + kernel_symbol_data(const base_type& _base) + : base_type{_base} + , formatted_kernel_name{tool::format_name(CHECK_NOTNULL(_base.kernel_name))} + , demangled_kernel_name{common::cxx_demangle(CHECK_NOTNULL(_base.kernel_name))} + , truncated_kernel_name{common::truncate_name(demangled_kernel_name)} {} + + std::string formatted_kernel_name = {}; + std::string demangled_kernel_name = {}; + std::string truncated_kernel_name = {}; }; -TraceBuffer hsa_api_buffer("HSA API", - 0x200000, - [](hsa_api_trace_entry_t* entry) { - TracerFlushRecord( - entry, - ROCPROFILER_CALLBACK_TRACING_HSA_API); - }); +using kernel_symbol_data_map_t = std::unordered_map; +auto kernel_data = common::Synchronized{}; +auto name_info = get_buffer_id_names(); -rocprofiler_tool_callback_name_info_t name_info; +auto& +get_client_ctx() +{ + static rocprofiler_context_id_t context_id; + return context_id; +} void -TracerFlushRecord(void* data, rocprofiler_callback_tracing_kind_t kind) +flush() { - if(kind == ROCPROFILER_CALLBACK_TRACING_HSA_API) + for(auto itr : {get_memory_copy_trace_buffer(), + get_kernel_trace_buffer(), + get_counter_collection_buffer(), + get_hsa_api_trace_buffer()}) { - auto* entry = reinterpret_cast(data); - get_hsa_api_file() << "ACTIVITY_DOMAIN_HSA_API" - << "," - << "\"" << entry->api_name << "\"" - << "," << entry->begin_timestamp << "," << entry->end_timestamp << "," - << entry->record.correlation_id.internal << '\n'; + if(itr.handle > 0) ROCPROFILER_CALL(rocprofiler_flush_buffer(itr), "buffer flush"); } } @@ -214,35 +208,11 @@ rocprofiler_tracing_callback(rocprofiler_callback_tracing_record_t record, rocprofiler_user_data_t* user_data, void* data) { - if(record.kind == ROCPROFILER_CALLBACK_TRACING_HIP_API) - { - // To be implemented - throw std::runtime_error{"not implemented"}; - } + throw std::runtime_error{"not implemented"}; - else if(record.kind == ROCPROFILER_CALLBACK_TRACING_HSA_API) - { - auto timestamp = rocprofiler_timestamp_t{}; - ROCPROFILER_CALL(rocprofiler_get_timestamp(×tamp), "timestamp failed"); - - if(record.phase == ROCPROFILER_CALLBACK_PHASE_ENTER) - { - user_data->value = timestamp; - } - else if(record.phase == ROCPROFILER_CALLBACK_PHASE_EXIT) - { - const auto* info_name_str = name_info.operation_names[record.kind][record.operation]; - hsa_api_trace_entry_t& entry = - hsa_api_buffer.Emplace(user_data->value, timestamp, record, info_name_str); - entry.valid.store(TRACE_ENTRY_COMPLETE, std::memory_order_release); - } - } - else if(record.kind == ROCPROFILER_CALLBACK_TRACING_MARKER_API) - { - // To be implemented - throw std::runtime_error{"not implemented"}; - } - (void) (data); + (void) record; + (void) user_data; + (void) data; } void @@ -255,38 +225,22 @@ code_object_tracing_callback(rocprofiler_callback_tracing_record_t record, { if(record.phase == ROCPROFILER_CALLBACK_PHASE_UNLOAD) { - // flush the buffer to ensure that any lookups for the client kernel names for the code - // object are completed - auto flush_status = rocprofiler_flush_buffer(tool_buffer); - if(flush_status != ROCPROFILER_STATUS_ERROR_BUFFER_BUSY) - ROCPROFILER_CALL(flush_status, "buffer flush"); + flush(); } } if(record.kind == ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT && record.operation == ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER) { - auto* sym_data = - static_cast( - record.payload); - + auto* sym_data = static_cast(record.payload); if(record.phase == ROCPROFILER_CALLBACK_PHASE_LOAD) { - auto kernel_name = - std::regex_replace(sym_data->kernel_name, std::regex{"(\\.kd)$"}, ""); - int demangle_status = 0; - kernel_name = cxa_demangle(kernel_name, &demangle_status); - std::unique_lock lock(kernel_data_mutex); - kernel_data.emplace(sym_data->kernel_id, kernel_name); + kernel_data.wlock( + [](kernel_symbol_data_map_t& kdata, rocprofiler_kernel_symbol_data_t* sym_data_v) { + kdata.emplace(sym_data_v->kernel_id, kernel_symbol_data{*sym_data_v}); + }, + sym_data); } - // The map entry cannot be erased here - // since we are tracing the kernel symbols here not the kernel dispatch - // else if(record.phase == ROCPROFILER_CALLBACK_PHASE_UNLOAD) - //{ - - // kernel_data.erase(data->kernel_id); - - //} } (void) user_data; @@ -294,174 +248,369 @@ code_object_tracing_callback(rocprofiler_callback_tracing_record_t record, } void -kernel_tracing_callback(rocprofiler_context_id_t context, - rocprofiler_buffer_id_t buffer_id, - rocprofiler_record_header_t** headers, - size_t num_headers, - void* user_data, - uint64_t /*drop_count*/) +buffered_callback(rocprofiler_context_id_t /*context*/, + rocprofiler_buffer_id_t /*buffer_id*/, + rocprofiler_record_header_t** headers, + size_t num_headers, + void* /*user_data*/, + uint64_t /*drop_count*/) { + static auto _sync = std::mutex{}; + auto _lk = std::lock_guard{_sync}; + if(num_headers == 0) - throw std::runtime_error{ - "rocprofiler invoked a buffer callback with no headers. this should never happen"}; + throw std::runtime_error{"rocprofiler invoked a buffer callback with no headers " + "this should never happen"}; + else if(headers == nullptr) throw std::runtime_error{"rocprofiler invoked a buffer callback with a null pointer to the " "array of headers. this should never happen"}; - auto kernel_trace_ss = std::stringstream{}; for(size_t i = 0; i < num_headers; ++i) { auto* header = headers[i]; - if(header == nullptr) + if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING) { - throw std::runtime_error{ - "rocprofiler provided a null pointer to header. this should never happen"}; - } - else if(header->hash != - rocprofiler_record_header_compute_hash(header->category, header->kind)) - { - throw std::runtime_error{"rocprofiler_record_header_t (category | kind) != hash"}; - } - else if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING && - header->kind == ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH) - { - auto* record = - static_cast(header->payload); - std::string kernel_name; + if(header->kind == ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH) { - std::shared_lock lock(kernel_data_mutex); - kernel_name = kernel_data.at(record->kernel_id); - } + auto* record = static_cast( + header->payload); + std::string kernel_name = kernel_data.rlock( + [](const kernel_symbol_data_map_t& kdata, + rocprofiler_buffer_tracing_kernel_dispatch_record_t* record_v) { + return kdata.at(record_v->kernel_id).formatted_kernel_name; + }, + record); - kernel_trace_ss << record->agent_id.handle << "," << record->queue_id.handle << "," - << record->kernel_id << ",\"" << kernel_name << "\"," << context.handle - << "," << buffer_id.handle << "," << record->correlation_id.internal - << "," << record->kind << "," << record->start_timestamp << "," - << record->end_timestamp << "," << record->private_segment_size << "," - << record->group_segment_size << "," << record->workgroup_size.x << "," - << record->workgroup_size.y << "," << record->workgroup_size.z << "," - << record->grid_size.x << "," << record->grid_size.y << "," - << record->grid_size.z << '\n'; + auto kernel_trace_ss = std::stringstream{}; + tool::csv::kernel_trace_csv_encoder::write_row( + kernel_trace_ss, + name_info.kind_names.at(record->kind), + record->agent_id.handle, + record->queue_id.handle, + record->kernel_id, + std::move(kernel_name), + record->correlation_id.internal, + record->start_timestamp, + record->end_timestamp, + record->private_segment_size, + record->group_segment_size, + record->workgroup_size.x, + record->workgroup_size.y, + record->workgroup_size.z, + record->grid_size.x, + record->grid_size.y, + record->grid_size.z); + + get_kernel_trace_file() << kernel_trace_ss.str(); + } + else if(header->kind == ROCPROFILER_BUFFER_TRACING_HSA_API) + { + auto* record = + static_cast(header->payload); + + auto hsa_trace_ss = std::stringstream{}; + tool::csv::hsa_csv_encoder::write_row( + hsa_trace_ss, + name_info.kind_names.at(record->kind), + name_info.operation_names.at(record->kind).at(record->operation), + getpid(), + record->thread_id, + record->correlation_id.internal, + record->start_timestamp, + record->end_timestamp); + + get_hsa_api_file() << hsa_trace_ss.str(); + } + else if(header->kind == ROCPROFILER_BUFFER_TRACING_MEMORY_COPY) + { + auto* record = + static_cast(header->payload); + + auto memory_copy_trace_ss = std::stringstream{}; + tool::csv::memory_copy_csv_encoder::write_row( + memory_copy_trace_ss, + name_info.kind_names.at(record->kind), + name_info.operation_names.at(record->kind).at(record->operation), + record->src_agent_id.handle, + record->dst_agent_id.handle, + record->correlation_id.internal, + record->start_timestamp, + record->end_timestamp); + + get_memory_copy_trace_file() << memory_copy_trace_ss.str(); + } + else + { + LOG(FATAL) << fmt::format( + "unsupported category + kind: {} + {}", header->category, header->kind); + } + } + + if(header->category == ROCPROFILER_BUFFER_CATEGORY_COUNTERS && header->kind == 0) + { + auto* profiler_record = static_cast(header->payload); + rocprofiler_tool_kernel_properties_t kernel_properties = + GetKernelProperties(profiler_record->corr_id.internal); + rocprofiler_counter_id_t counter_id; + const char* counter_name; + size_t size, pos; + rocprofiler_query_record_counter_id(profiler_record->id, &counter_id); + rocprofiler_query_counter_name(counter_id, &counter_name, &size); + rocprofiler_query_record_dimension_position(profiler_record->id, 0, &pos); + + auto counter_collection_ss = std::stringstream{}; + counter_collection_ss << counter_id.handle << "," + << kernel_properties.gpu_agent.id.handle << "," + << kernel_properties.queue_id.handle << "," << getpid() << "," + << kernel_properties.thread_id << ","; + counter_collection_ss << kernel_properties.grid_size << "," + << kernel_properties.kernel_name << "," + << kernel_properties.workgroup_size << "," + << ((kernel_properties.lds_size + (lds_block_size - 1)) & + ~(lds_block_size - 1)) + << "," << kernel_properties.scratch_size << "," + << kernel_properties.arch_vgpr_count << "," + << kernel_properties.sgpr_count << ","; + /* + Iterate through the N dimensional that is obtained for the counter. + given instance id what is the counter id + given counter id what is the counter name + given instance how many dimension + iterate through dimensions + what is the dimension id + what is the dimension name + what pos in the dimension. + */ + // ss << counter_name << "[" << info.name << "," << pos << "]" << ","; + // ss << profiler_record->counter_value << "\n"; + counter_collection_ss << counter_name << "[" + << "," << pos << "]" + << ","; + counter_collection_ss << counter_name << ","; + counter_collection_ss << profiler_record->counter_value << "\n"; + get_counter_collection_file() << counter_collection_ss.str() << "\n"; } } - - static auto _sync = std::mutex{}; - auto _lk = std::unique_lock{_sync}; - if(get_kernel_trace_file()) - get_kernel_trace_file() << kernel_trace_ss.str(); - else - std::cerr << "kernel trace file already closed: " << kernel_trace_ss.str(); - - (void) (user_data); } -rocprofiler_tool_callback_name_info_t -get_callback_id_names() +using counter_vec_t = std::vector; +using agent_counter_map_t = + std::unordered_map>; + +// this function creates a rocprofiler profile config on the first entry +auto +get_agent_profile(const rocprofiler_agent_t* agent) { - auto cb_name_info = rocprofiler_tool_callback_name_info_t{}; - // - // callback for each kind operation - // - static auto tracing_kind_operation_cb = - [](rocprofiler_callback_tracing_kind_t kindv, uint32_t operation, void* data_v) { - auto* name_info_v = static_cast(data_v); + static auto data = common::Synchronized{}; - if(kindv == ROCPROFILER_CALLBACK_TRACING_HSA_API) + auto profile = std::optional{}; + data.ulock( + [agent, &profile](const agent_counter_map_t& data_v) { + auto itr = data_v.find(agent); + if(itr != data_v.end()) { - const char* name = nullptr; - ROCPROFILER_CALL(rocprofiler_query_callback_tracing_kind_operation_name( - kindv, operation, &name, nullptr), - "query callback failed"); - if(name) name_info_v->operation_names[kindv][operation] = name; + profile = itr->second; + return true; + } + return false; + }, + [agent, &profile](agent_counter_map_t& data_v) { + auto counters_v = counter_vec_t{}; + ROCPROFILER_CALL( + rocprofiler_iterate_agent_supported_counters( + *agent, + [](rocprofiler_counter_id_t* counters, size_t num_counters, void* user_data) { + auto* vec = static_cast(user_data); + for(size_t i = 0; i < num_counters; i++) + { + const char* name = nullptr; + size_t len = 0; + + ROCPROFILER_CALL( + rocprofiler_query_counter_name(counters[i], &name, &len), + "Could not query name"); + + if(name && len > 0) + { + if(tool::get_config().counters.count(name) > 0) + vec->emplace_back(counters[i]); + } + } + return ROCPROFILER_STATUS_SUCCESS; + }, + static_cast(&counters_v)), + "iterate agent supported counters"); + + if(!counters_v.empty()) + { + auto profile_v = rocprofiler_profile_config_id_t{}; + ROCPROFILER_CALL(rocprofiler_create_profile_config( + *agent, counters_v.data(), counters_v.size(), &profile_v), + "Could not construct profile cfg"); + profile = profile_v; } - if(kindv == ROCPROFILER_CALLBACK_TRACING_HIP_API) - { - const char* name = nullptr; - ROCPROFILER_CALL(rocprofiler_query_callback_tracing_kind_operation_name( - kindv, operation, &name, nullptr), - "query callback failed"); - if(name) name_info_v->operation_names[kindv][operation] = name; - } - return 0; - }; + data_v.emplace(agent, profile); + return true; + }); - // - // callback for each callback kind (i.e. domain) - // - static auto tracing_kind_cb = [](rocprofiler_callback_tracing_kind_t kind, void* data) { - // store the callback kind name - auto* name_info_v = static_cast(data); - const char* name = nullptr; - ROCPROFILER_CALL(rocprofiler_query_callback_tracing_kind_name(kind, &name, nullptr), - "query callback failed"); + return profile; +} - if(name) name_info_v->kind_names[kind] = name; +void +dispatch_callback(rocprofiler_queue_id_t queue_id, + const rocprofiler_agent_t* agent, + rocprofiler_correlation_id_t correlation_id, + const hsa_kernel_dispatch_packet_t* dispatch_packet, + uint64_t kernel_id, + void* /*callback_data_args*/, + rocprofiler_profile_config_id_t* config) +{ + rocprofiler_tool_kernel_properties_t kernel_properties; + const auto& kernel_info = + kernel_data.rlock([](const kernel_symbol_data_map_t& kdata, + uint64_t kernel_id_v) { return kdata.at(kernel_id_v); }, + kernel_id); - if(kind == ROCPROFILER_CALLBACK_TRACING_HSA_API) + auto is_targeted_kernel = [&kernel_info]() { + for(const auto& name : tool::get_config().kernel_names) { - ROCPROFILER_CALL(rocprofiler_iterate_callback_tracing_kind_operations( - kind, tracing_kind_operation_cb, static_cast(data)), - "query callback failed"); + if(name == kernel_info.truncated_kernel_name) + return true; + else + { + auto dkernel_name = std::string_view{kernel_info.demangled_kernel_name}; + auto pos = dkernel_name.find(name); + // if the demangled kernel name contains name and the next character is '(' then + // mark as found + if(pos != std::string::npos && (pos + 1) < dkernel_name.size() && + dkernel_name.at(pos + 1) == '(') + return true; + } } - return 0; + return false; }; - ROCPROFILER_CALL(rocprofiler_iterate_callback_tracing_kinds(tracing_kind_cb, - static_cast(&cb_name_info)), - "iterate_callback failed"); + if(!is_targeted_kernel()) return; - return cb_name_info; + auto profile = get_agent_profile(agent); + + if(profile) + { + kernel_properties.kernel_name = kernel_info.formatted_kernel_name; + kernel_properties.queue_id = queue_id; + kernel_properties.gpu_agent = *agent; + kernel_properties.thread_id = common::get_tid(); + populate_kernel_properties_data(&kernel_properties, dispatch_packet); + SetKernelProperties(correlation_id.internal, kernel_properties); + + *config = *profile; + } } +rocprofiler_client_finalize_t client_finalizer = nullptr; +rocprofiler_client_id_t* client_identifier = nullptr; + int -tool_init(rocprofiler_client_finalize_t /*fini_func*/, void* tool_data) +tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) { - name_info = get_callback_id_names(); + client_finalizer = fini_func; - ROCPROFILER_CALL(rocprofiler_create_context(&context_id), "create context failed"); + ROCPROFILER_CALL(rocprofiler_create_context(&get_client_ctx()), "create context failed"); - if(common::get_env("ROCPROFILER_KERNEL_TRACE", false)) + ROCPROFILER_CALL( + rocprofiler_configure_callback_tracing_service(get_client_ctx(), + ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT, + nullptr, + 0, + code_object_tracing_callback, + nullptr), + "code object tracing configure failed"); + + if(tool::get_config().kernel_trace) { - ROCPROFILER_CALL( - rocprofiler_configure_callback_tracing_service(context_id, - ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT, - nullptr, - 0, - code_object_tracing_callback, - nullptr), - "tracing configure failed"); - - ROCPROFILER_CALL(rocprofiler_create_buffer(context_id, + ROCPROFILER_CALL(rocprofiler_create_buffer(get_client_ctx(), 4096, 2048, ROCPROFILER_BUFFER_POLICY_LOSSLESS, - kernel_tracing_callback, + buffered_callback, tool_data, - &tool_buffer), + &get_kernel_trace_buffer()), "buffer creation"); ROCPROFILER_CALL( - rocprofiler_configure_buffer_tracing_service( - context_id, ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, nullptr, 0, tool_buffer), + rocprofiler_configure_buffer_tracing_service(get_client_ctx(), + ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, + nullptr, + 0, + get_kernel_trace_buffer()), "buffer tracing service for kernel dispatch configure"); } - if(common::get_env("ROCPROFILER_HSA_API_TRACE", false)) + if(tool::get_config().memory_copy_trace) { - // Requesting all operations + ROCPROFILER_CALL(rocprofiler_create_buffer(get_client_ctx(), + 4096, + 2048, + ROCPROFILER_BUFFER_POLICY_LOSSLESS, + buffered_callback, + nullptr, + &get_memory_copy_trace_buffer()), + "create memory copy buffer"); + ROCPROFILER_CALL( - rocprofiler_configure_callback_tracing_service(context_id, - ROCPROFILER_CALLBACK_TRACING_HSA_API, - nullptr, - 0, - rocprofiler_tracing_callback, - nullptr), - "tracing configure failed"); + rocprofiler_configure_buffer_tracing_service(get_client_ctx(), + ROCPROFILER_BUFFER_TRACING_MEMORY_COPY, + nullptr, + 0, + get_memory_copy_trace_buffer()), + "buffer tracing service for memory copy configure"); } - ROCPROFILER_CALL(rocprofiler_start_context(context_id), "start context failed"); + if(tool::get_config().hsa_api_trace) + { + ROCPROFILER_CALL(rocprofiler_create_buffer(get_client_ctx(), + 4096, + 2048, + ROCPROFILER_BUFFER_POLICY_LOSSLESS, + buffered_callback, + tool_data, + &get_hsa_api_trace_buffer()), + "buffer creation"); + + ROCPROFILER_CALL( + rocprofiler_configure_buffer_tracing_service(get_client_ctx(), + ROCPROFILER_BUFFER_TRACING_HSA_API, + nullptr, + 0, + get_hsa_api_trace_buffer()), + "buffer tracing service for memory copy configure"); + } + + if(tool::get_config().counter_collection) + { + ROCPROFILER_CALL(rocprofiler_create_buffer(get_client_ctx(), + 4096, + 2048, + ROCPROFILER_BUFFER_POLICY_LOSSLESS, + buffered_callback, + nullptr, + &get_counter_collection_buffer()), + "buffer creation failed"); + + ROCPROFILER_CALL( + rocprofiler_configure_buffered_dispatch_profile_counting_service( + get_client_ctx(), get_counter_collection_buffer(), dispatch_callback, nullptr), + "Could not setup buffered service"); + } + + ROCPROFILER_CALL(rocprofiler_start_context(get_client_ctx()), "start context failed"); + + std::atexit([]() { + if(client_finalizer && client_identifier) client_finalizer(*client_identifier); + }); return 0; } @@ -469,8 +618,12 @@ tool_init(rocprofiler_client_finalize_t /*fini_func*/, void* tool_data) void tool_fini(void* tool_data) { - rocprofiler_flush_buffer(tool_buffer); - rocprofiler_stop_context(context_id); + client_identifier = nullptr; + client_finalizer = nullptr; + + flush(); + rocprofiler_stop_context(get_client_ctx()); + (void) (tool_data); } @@ -480,14 +633,17 @@ rocprofiler_configure(uint32_t /*version*/, uint32_t priority, rocprofiler_client_id_t* id) { + common::init_logging("ROCPROF_LOG_LEVEL"); + FLAGS_colorlogtostderr = true; + // only activate if main tool if(priority > 0) return nullptr; // set the client name - id->name = "rocporfiler-tool"; + id->name = "rocprofiler-tool"; // store client info - // client::client_id = id; + client_identifier = id; // create configure data static auto cfg = rocprofiler_tool_configure_result_t{ diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.cpp index 62282d9400..7c41083060 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.cpp @@ -26,6 +26,7 @@ #include #include "lib/common/container/stable_vector.hpp" +#include "lib/common/static_object.hpp" #include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/rocprofiler-sdk/context/domain.hpp" @@ -70,24 +71,26 @@ get_buffer_offset() bool is_valid_buffer_id(rocprofiler_buffer_id_t id) { - auto nbuffers = get_buffers().size(); + if(!get_buffers()) return false; + auto nbuffers = get_buffers()->size(); auto offset = get_buffer_offset(); return (id.handle >= offset && id.handle < (offset + nbuffers)); } -unique_buffer_vec_t& +unique_buffer_vec_t* get_buffers() { - static auto _v = unique_buffer_vec_t{reserve_size_t{unique_buffer_vec_t::chunk_size}}; + static auto*& _v = common::static_object::construct( + reserve_size_t{unique_buffer_vec_t::chunk_size}); return _v; } instance* get_buffer(rocprofiler_buffer_id_t buffer_id) { - if(is_valid_buffer_id(buffer_id)) + if(is_valid_buffer_id(buffer_id) && get_buffers()) { - for(auto& itr : get_buffers()) + for(auto& itr : *get_buffers()) { if(itr && itr->buffer_id == buffer_id.handle) { @@ -101,6 +104,8 @@ get_buffer(rocprofiler_buffer_id_t buffer_id) std::optional allocate_buffer() { + if(registration::get_fini_status() > 0) return std::nullopt; + // ensure buffer has thread to handle flushing it static auto _init_threads_once = std::once_flag{}; std::call_once(_init_threads_once, []() { internal_threading::initialize(); }); @@ -109,13 +114,13 @@ allocate_buffer() auto _lk = std::unique_lock{get_buffers_mutex()}; // initial context identifier number - auto _idx = get_buffer_offset() + get_buffers().size(); + auto _idx = get_buffer_offset() + CHECK_NOTNULL(get_buffers())->size(); // make space in registered - get_buffers().emplace_back(nullptr); + CHECK_NOTNULL(get_buffers())->emplace_back(nullptr); // create an entry in the registered - auto& _cfg_v = get_buffers().back(); + auto& _cfg_v = CHECK_NOTNULL(get_buffers())->back(); _cfg_v = allocator::make_unique_static(); auto* _cfg = _cfg_v.get(); @@ -157,7 +162,7 @@ flush(rocprofiler_buffer_id_t buffer_id, bool wait) auto idx = buff->buffer_idx++; auto _task = [buffer_id, idx, offset]() { - auto& buff_v = get_buffers().at(buffer_id.handle - offset); + auto& buff_v = CHECK_NOTNULL(get_buffers())->at(buffer_id.handle - offset); auto& buff_internal_v = buff_v->get_internal_buffer(idx); if(!buff_internal_v.is_empty()) @@ -232,8 +237,8 @@ rocprofiler_create_buffer(rocprofiler_context_id_t context, if(!opt_buff_id) return ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND; buffer_id->handle = opt_buff_id->handle; - auto& buff = rocprofiler::buffer::get_buffers().at(opt_buff_id->handle - - rocprofiler::buffer::get_buffer_offset()); + auto& buff = CHECK_NOTNULL(rocprofiler::buffer::get_buffers()) + ->at(opt_buff_id->handle - rocprofiler::buffer::get_buffer_offset()); // allocate the buffers. if it is lossless, we allocate a second buffer to store data while // other buffer is being flushed @@ -264,8 +269,8 @@ rocprofiler_destroy_buffer(rocprofiler_buffer_id_t buffer_id) return ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND; auto offset = rocprofiler::buffer::get_buffer_offset(); - auto& buffers = rocprofiler::buffer::get_buffers(); - auto& buff = buffers.at(buffer_id.handle - offset); + auto* buffers = CHECK_NOTNULL(rocprofiler::buffer::get_buffers()); + auto& buff = buffers->at(buffer_id.handle - offset); // buffer is currently being flushed or destroyed if(buff->syncer.test_and_set()) return ROCPROFILER_STATUS_ERROR_BUFFER_BUSY; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.hpp index 512f469ba8..794213a2fa 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.hpp @@ -71,7 +71,7 @@ is_valid_buffer_id(rocprofiler_buffer_id_t id); std::optional allocate_buffer(); -unique_buffer_vec_t& +unique_buffer_vec_t* get_buffers(); instance* diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer_tracing.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer_tracing.cpp index 063ca8db49..88da2571d3 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer_tracing.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer_tracing.cpp @@ -25,6 +25,7 @@ #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/rocprofiler-sdk/context/domain.hpp" +#include "lib/rocprofiler-sdk/hsa/async_copy.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" #include "lib/rocprofiler-sdk/marker/marker.hpp" #include "lib/rocprofiler-sdk/registration.hpp" @@ -96,6 +97,8 @@ rocprofiler_configure_buffer_tracing_service(rocprofiler_context_id_t c if(!ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; + if(buffer_id.handle == 0) return ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND; + constexpr auto invalid_buffer_id = rocprofiler_buffer_id_t{std::numeric_limits::max()}; @@ -147,6 +150,8 @@ rocprofiler_query_buffer_tracing_kind_operation_name(rocprofiler_buffer_tracing_ const char* val = nullptr; if(kind == ROCPROFILER_BUFFER_TRACING_HSA_API) val = rocprofiler::hsa::name_by_id(operation); + else if(kind == ROCPROFILER_BUFFER_TRACING_MEMORY_COPY) + val = rocprofiler::hsa::async_copy::name_by_id(operation); else if(kind == ROCPROFILER_BUFFER_TRACING_MARKER_API) val = rocprofiler::marker::name_by_id(operation); else @@ -187,6 +192,8 @@ rocprofiler_iterate_buffer_tracing_kind_operations( auto ops = std::vector{}; if(kind == ROCPROFILER_BUFFER_TRACING_HSA_API) ops = rocprofiler::hsa::get_ids(); + else if(kind == ROCPROFILER_BUFFER_TRACING_MEMORY_COPY) + ops = rocprofiler::hsa::async_copy::get_ids(); else if(kind == ROCPROFILER_BUFFER_TRACING_MARKER_API) ops = rocprofiler::marker::get_ids(); else diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp index ae63ab60e8..415c8d2c72 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp @@ -420,11 +420,11 @@ deregister_client_contexts(rocprofiler_client_id_t client_id) { for(auto& itr : get_registered_contexts_impl()) { - if(itr->client_idx == client_id.handle) + if(itr->client_idx == client_id.handle && buffer::get_buffers()) { - for(auto& bitr : buffer::get_buffers()) + for(auto& bitr : *buffer::get_buffers()) { - if(bitr->context_id == itr->context_idx) bitr.reset(); + if(bitr && bitr->context_id == itr->context_idx) bitr.reset(); } itr.reset(); } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp index 6714c4dc52..28df0cfcb8 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp @@ -23,6 +23,7 @@ #include "lib/rocprofiler-sdk/hsa/async_copy.hpp" #include "lib/common/defines.hpp" #include "lib/common/scope_destructor.hpp" +#include "lib/common/static_object.hpp" #include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/agent.hpp" #include "lib/rocprofiler-sdk/buffer.hpp" @@ -61,12 +62,78 @@ namespace rocprofiler { namespace hsa { +namespace async_copy +{ namespace { using context_t = context::context; using context_array_t = common::container::small_vector; using external_corr_id_map_t = std::unordered_map; +template +struct async_copy_info; + +#define SPECIALIZE_ASYNC_COPY_INFO(DIRECTION) \ + template <> \ + struct async_copy_info \ + { \ + static constexpr auto operation_idx = ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_##DIRECTION; \ + static constexpr auto name = #DIRECTION; \ + }; + +SPECIALIZE_ASYNC_COPY_INFO(NONE) +SPECIALIZE_ASYNC_COPY_INFO(HOST_TO_HOST) +SPECIALIZE_ASYNC_COPY_INFO(HOST_TO_DEVICE) +SPECIALIZE_ASYNC_COPY_INFO(DEVICE_TO_HOST) +SPECIALIZE_ASYNC_COPY_INFO(DEVICE_TO_DEVICE) + +#undef SPECIALIZE_ASYNC_COPY_INFO + +template +const char* +name_by_id(const uint32_t id, std::index_sequence) +{ + if(Idx == id) return async_copy_info::name; + if constexpr(sizeof...(IdxTail) > 0) + return name_by_id(id, std::index_sequence{}); + else + return nullptr; +} + +template +uint32_t +id_by_name(const char* name, std::index_sequence) +{ + if(std::string_view{async_copy_info::name} == std::string_view{name}) + return async_copy_info::operation_idx; + if constexpr(sizeof...(IdxTail) > 0) + return id_by_name(name, std::index_sequence{}); + else + return ROCPROFILER_HSA_API_ID_NONE; +} + +template +void +get_ids(std::vector& _id_list, std::index_sequence) +{ + auto _emplace = [](auto& _vec, uint32_t _v) { + if(_v < ROCPROFILER_HSA_API_ID_LAST) _vec.emplace_back(_v); + }; + + (_emplace(_id_list, async_copy_info::operation_idx), ...); +} + +template +void +get_names(std::vector& _name_list, std::index_sequence) +{ + auto _emplace = [](auto& _vec, const char* _v) { + if(_v != nullptr && strnlen(_v, 1) > 0) _vec.emplace_back(_v); + }; + + (_emplace(_name_list, async_copy_info::name), ...); +} + bool context_filter(const context::context* ctx) { @@ -90,6 +157,13 @@ struct async_copy_data external_corr_id_map_t extern_corr_ids = {}; }; +auto* +get_active_signals() +{ + static auto* _v = common::static_object>::construct(); + return _v; +} + template constexpr Tp* convert_hsa_handle(Up _hsa_object) @@ -197,6 +271,7 @@ async_copy_handler(hsa_signal_value_t signal_value, void* arg) { ROCP_HSA_TABLE_CALL(ERROR, get_table().core_->hsa_signal_destroy_fn(_data->rocp_signal)); delete _data; + get_active_signals()->fetch_sub(1); } return (signal_value > 0); @@ -367,6 +442,10 @@ async_copy_impl(Args... args) return invoke( get_next_dispatch(), std::move(_tied_args), std::make_index_sequence{}); } + else + { + get_active_signals()->fetch_add(1); + } } { @@ -384,6 +463,8 @@ async_copy_impl(Args... args) ROCP_HSA_TABLE_CALL(ERROR, get_table().core_->hsa_signal_destroy_fn(_data->rocp_signal)) << ":: failed to destroy signal after async handler failed"; + get_active_signals()->fetch_sub(1); + delete _data; return invoke( get_next_dispatch(), std::move(_tied_args), std::make_index_sequence{}); @@ -455,20 +536,64 @@ using async_copy_index_seq_t = std::index_sequence; } // namespace +// check out the assembly here... this compiles to a switch statement +const char* +name_by_id(uint32_t id) +{ + return name_by_id(id, std::make_index_sequence{}); +} + +uint32_t +id_by_name(const char* name) +{ + return id_by_name(name, + std::make_index_sequence{}); +} + +std::vector +get_ids() +{ + auto _data = std::vector{}; + _data.reserve(ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST); + get_ids(_data, std::make_index_sequence{}); + return _data; +} + +std::vector +get_names() +{ + auto _data = std::vector{}; + _data.reserve(ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST); + get_names(_data, std::make_index_sequence{}); + return _data; +} +} // namespace async_copy + void async_copy_init(hsa_api_table_t* _orig) { if(_orig) { - async_copy_save(_orig, async_copy_index_seq_t{}); + async_copy::async_copy_save(_orig, async_copy::async_copy_index_seq_t{}); - auto ctxs = context::get_registered_contexts(context_filter); + auto ctxs = context::get_registered_contexts(async_copy::context_filter); if(!ctxs.empty()) { _orig->amd_ext_->hsa_amd_profiling_async_copy_enable_fn(true); - async_copy_wrap(_orig, async_copy_index_seq_t{}); + async_copy::async_copy_wrap(_orig, async_copy::async_copy_index_seq_t{}); } } } + +void +async_copy_fini() +{ + if(!async_copy::get_active_signals()) return; + while(async_copy::get_active_signals()->load() > 0) + { + std::this_thread::yield(); + std::this_thread::sleep_for(std::chrono::milliseconds{50}); + } +} } // namespace hsa } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.hpp index a99071c351..1324f9b107 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.hpp @@ -30,7 +30,25 @@ namespace rocprofiler { namespace hsa { +namespace async_copy +{ +const char* +name_by_id(uint32_t id); + +uint32_t +id_by_name(const char* name); + +std::vector +get_names(); + +std::vector +get_ids(); +} // namespace async_copy + void async_copy_init(hsa_api_table_t* _orig); + +void +async_copy_fini(); } // namespace hsa } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/registration.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/registration.cpp index 61e61887b6..ac518d1960 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/registration.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/registration.cpp @@ -557,6 +557,7 @@ finalize() static auto _once = std::once_flag{}; std::call_once(_once, []() { set_fini_status(-1); + hsa::async_copy_fini(); hsa::code_object_shutdown(); if(get_init_status() > 0) { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/buffer.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/buffer.cpp index 07b3f41c82..5fcc2a5197 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/buffer.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/buffer.cpp @@ -40,14 +40,14 @@ TEST(rocprofiler_lib, buffer) namespace buffer = ::rocprofiler::buffer; namespace common = ::rocprofiler::common; - ASSERT_EQ(buffer::get_buffers().size(), 0) + ASSERT_EQ(buffer::get_buffers()->size(), 0) << "no buffers should have been created at this point"; auto buffer_id = buffer::allocate_buffer(); EXPECT_TRUE(buffer_id) << "failed to allocate buffer"; EXPECT_GT(buffer_id->handle, 0); EXPECT_TRUE(buffer::is_valid_buffer_id(*buffer_id)) << "id=" << buffer_id->handle; - ASSERT_EQ(buffer::get_buffers().size(), 1) << "incorrect number of buffers created"; + ASSERT_EQ(buffer::get_buffers()->size(), 1) << "incorrect number of buffers created"; // get pointer to buffer auto* buffer_v = buffer::get_buffer(*buffer_id); diff --git a/projects/rocprofiler-sdk/tests/CMakeLists.txt b/projects/rocprofiler-sdk/tests/CMakeLists.txt index 97bc24170f..51b9ba203c 100644 --- a/projects/rocprofiler-sdk/tests/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/CMakeLists.txt @@ -42,3 +42,6 @@ add_subdirectory(apps) # validation tests add_subdirectory(kernel-tracing) add_subdirectory(async-copy-tracing) + +# rocprofv3 validation tests +add_subdirectory(rocprofv3) diff --git a/projects/rocprofiler-sdk/tests/apps/CMakeLists.txt b/projects/rocprofiler-sdk/tests/apps/CMakeLists.txt index c5ef92a52c..86d833ceff 100644 --- a/projects/rocprofiler-sdk/tests/apps/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/apps/CMakeLists.txt @@ -5,10 +5,15 @@ cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) project(rocprofiler-test-apps LANGUAGES C CXX) +set(CMAKE_BUILD_RPATH "\$ORIGIN:\$ORIGIN/../lib") + +# applications used by integration tests which DO NOT link to rocprofiler-sdk-roctx +add_subdirectory(simple-transpose) + set(CMAKE_BUILD_RPATH "\$ORIGIN:\$ORIGIN/../lib:$" ) -# applications used by integration tests +# applications used by integration tests which DO link to rocprofiler-sdk-roctx add_subdirectory(reproducible-runtime) add_subdirectory(transpose) diff --git a/projects/rocprofiler-sdk/tests/apps/simple-transpose/CMakeLists.txt b/projects/rocprofiler-sdk/tests/apps/simple-transpose/CMakeLists.txt new file mode 100644 index 0000000000..3310175e9d --- /dev/null +++ b/projects/rocprofiler-sdk/tests/apps/simple-transpose/CMakeLists.txt @@ -0,0 +1,47 @@ +# +# +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +if(NOT CMAKE_HIP_COMPILER) + find_program( + amdclangpp_EXECUTABLE + NAMES amdclang++ + HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm + PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm + PATH_SUFFIXES bin llvm/bin NO_CACHE) + mark_as_advanced(amdclangpp_EXECUTABLE) + + if(amdclangpp_EXECUTABLE) + set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}") + endif() +endif() + +project(rocprofiler-tool-test-app-transpose LANGUAGES CXX HIP) + +foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO) + if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "") + set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}") + endif() +endforeach() + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_EXTENSIONS OFF) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_HIP_STANDARD 17) +set(CMAKE_HIP_EXTENSIONS OFF) +set(CMAKE_HIP_STANDARD_REQUIRED ON) + +set_source_files_properties(simple-transpose.cpp PROPERTIES LANGUAGE HIP) +add_executable(simple-transpose) +target_sources(simple-transpose PRIVATE simple-transpose.cpp) +target_compile_options(simple-transpose PRIVATE -W -Wall -Wextra -Wpedantic -Wshadow + -Werror) + +find_package(Threads REQUIRED) +target_link_libraries(simple-transpose PRIVATE Threads::Threads) + +install( + TARGETS simple-transpose + DESTINATION bin + COMPONENT tests) diff --git a/projects/rocprofiler-sdk/tests/apps/simple-transpose/simple-transpose.cpp b/projects/rocprofiler-sdk/tests/apps/simple-transpose/simple-transpose.cpp new file mode 100644 index 0000000000..c7daa0f801 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/apps/simple-transpose/simple-transpose.cpp @@ -0,0 +1,162 @@ +/* +Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include + +// hip header file +#include "hip/hip_runtime.h" + +#define WIDTH 1024 + +#define NUM (WIDTH * WIDTH) + +#define THREADS_PER_BLOCK_X 4 +#define THREADS_PER_BLOCK_Y 4 +#define THREADS_PER_BLOCK_Z 1 + +#define HIP_API_CALL(CALL) \ + { \ + hipError_t error_ = (CALL); \ + if(error_ != hipSuccess) \ + { \ + auto _hip_api_print_lk = auto_lock_t{print_lock}; \ + fprintf(stderr, \ + "%s:%d :: HIP error : %s\n", \ + __FILE__, \ + __LINE__, \ + hipGetErrorString(error_)); \ + throw std::runtime_error("hip_api_call"); \ + } \ + } + +namespace +{ +using auto_lock_t = std::unique_lock; +auto print_lock = std::mutex{}; +} // namespace + +// Device (Kernel) function, it must be void +__global__ void +matrixTranspose(float* out, float* in, const int width) +{ + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; + + out[y * width + x] = in[x * width + y]; +} + +// CPU implementation of matrix transpose +void +matrixTransposeCPUReference(float* output, float* input, const unsigned int width) +{ + for(unsigned int j = 0; j < width; j++) + { + for(unsigned int i = 0; i < width; i++) + { + output[i * width + j] = input[j * width + i]; + } + } +} + +int +main() +{ + float* Matrix; + float* TransposeMatrix; + float* cpuTransposeMatrix; + + float* gpuMatrix; + float* gpuTransposeMatrix; + + hipDeviceProp_t devProp; + HIP_API_CALL(hipGetDeviceProperties(&devProp, 0)); + + std::cout << "Device name " << devProp.name << std::endl; + + int i; + int errors; + + Matrix = (float*) malloc(NUM * sizeof(float)); + TransposeMatrix = (float*) malloc(NUM * sizeof(float)); + cpuTransposeMatrix = (float*) malloc(NUM * sizeof(float)); + + // initialize the input data + for(i = 0; i < NUM; i++) + { + Matrix[i] = (float) i * 10.0f; + } + + // allocate the memory on the device side + HIP_API_CALL(hipMalloc((void**) &gpuMatrix, NUM * sizeof(float))); + HIP_API_CALL(hipMalloc((void**) &gpuTransposeMatrix, NUM * sizeof(float))); + + // Memory transfer from host to device + HIP_API_CALL(hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice)); + + // Lauching kernel from host + hipLaunchKernelGGL(matrixTranspose, + dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), + 0, + 0, + gpuTransposeMatrix, + gpuMatrix, + WIDTH); + + // Memory transfer from device to host + HIP_API_CALL( + hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost)); + + // CPU MatrixTranspose computation + matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); + + // verify the results + errors = 0; + double eps = 1.0E-6; + for(i = 0; i < NUM; i++) + { + if(std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) + { + errors++; + } + } + if(errors != 0) + { + printf("FAILED: %d errors\n", errors); + } + else + { + printf("PASSED!\n"); + } + + // free the resources on device side + HIP_API_CALL(hipFree(gpuMatrix)); + HIP_API_CALL(hipFree(gpuTransposeMatrix)); + + // free the resources on host side + free(Matrix); + free(TransposeMatrix); + free(cpuTransposeMatrix); + + return errors; +} diff --git a/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py b/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py index 0838b73eb2..1e03d16b03 100644 --- a/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py +++ b/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py @@ -154,26 +154,26 @@ def test_async_copy_direction(input_data): sdk_data = data["rocprofiler-sdk-json-tool"] # Direction values: - # -1 == ??? (unknown) - # 0 == H2H (host to host) - # 1 == H2D (host to device) - # 2 == D2H (device to host) - # 3 == D2D (device to device) - async_dir_cnt = dict([(idx, 0) for idx in range(-1, 4)]) + # 0 == ??? (unknown) + # 1 == H2H (host to host) + # 2 == H2D (host to device) + # 3 == D2H (device to host) + # 4 == D2D (device to device) + async_dir_cnt = dict([(idx, 0) for idx in range(0, 5)]) for itr in sdk_data["buffer_records"]["memory_copies"]: op_id = itr["operation"] - assert op_id > 0 - assert op_id < 3 + assert op_id > 1 + assert op_id < 4 async_dir_cnt[op_id] += 1 # in the transpose test which generates the input file, # two threads and the main thread (so three threads total) # each perform one H2D + one D2H memory copy - assert async_dir_cnt[-1] == 0 assert async_dir_cnt[0] == 0 - assert async_dir_cnt[1] == 6 + assert async_dir_cnt[1] == 0 assert async_dir_cnt[2] == 6 - assert async_dir_cnt[3] == 0 + assert async_dir_cnt[3] == 6 + assert async_dir_cnt[4] == 0 if __name__ == "__main__": diff --git a/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py b/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py index 58368bbf54..78f78392b9 100644 --- a/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py +++ b/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py @@ -156,23 +156,23 @@ def test_async_copy_direction(input_data): sdk_data = data["rocprofiler-sdk-json-tool"] # Direction values: - # -1 == ??? (unknown) - # 0 == H2H (host to host) - # 1 == H2D (host to device) - # 2 == D2H (device to host) - # 3 == D2D (device to device) - async_dir_cnt = dict([(idx, 0) for idx in range(-1, 4)]) + # 0 == ??? (unknown) + # 1 == H2H (host to host) + # 2 == H2D (host to device) + # 3 == D2H (device to host) + # 4 == D2D (device to device) + async_dir_cnt = dict([(idx, 0) for idx in range(0, 5)]) for itr in sdk_data["buffer_records"]["memory_copies"]: op_id = itr["operation"] async_dir_cnt[op_id] += 1 # in the reproducible-runtime test which generates the input file, # we don't expect any async memory copy operations - assert async_dir_cnt[-1] == 0 assert async_dir_cnt[0] == 0 assert async_dir_cnt[1] == 0 assert async_dir_cnt[2] == 0 assert async_dir_cnt[3] == 0 + assert async_dir_cnt[4] == 0 if __name__ == "__main__": diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt new file mode 100644 index 0000000000..8469100552 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt @@ -0,0 +1,43 @@ +# +# +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +project(rocprofiler-rocprofv3-tests LANGUAGES C CXX) + +set(CMAKE_INSTALL_DEFAULT_COMPONENT_NAME "tests") + +if(NOT CMAKE_BUILD_TYPE) + set(CMAKE_BUILD_TYPE + "Release" + CACHE STRING "" FORCE) +endif() + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_EXTENSIONS OFF) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_HIP_STANDARD 17) +set(CMAKE_HIP_EXTENSIONS OFF) +set(CMAKE_HIP_STANDARD_REQUIRED ON) + +enable_testing() +include(CTest) + +# needed for validation +find_package(Python3 REQUIRED) + +execute_process( + COMMAND ${Python3_EXECUTABLE} -c "import pandas" + RESULT_VARIABLE PYTHON_PANDAS_IMPORT + OUTPUT_QUIET) + +if(NOT ${PYTHON_PANDAS_IMPORT} EQUAL 0) + message( + FATAL_ERROR + "The 'pandas' Python3 package is not installed. Please install it using the following command: '${Python3_EXECUTABLE} -m pip install pandas'" + ) +endif() + +add_subdirectory(tracing) +add_subdirectory(counter-collection) +add_subdirectory(tracing-plus-cc) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/CMakeLists.txt new file mode 100644 index 0000000000..36ee2980d2 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/CMakeLists.txt @@ -0,0 +1,7 @@ +foreach(FILENAME conftest.py pytest.ini) + configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} + ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY) +endforeach() + +add_subdirectory(input1) +add_subdirectory(input2) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/conftest.py new file mode 100644 index 0000000000..f7ed55012e --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/conftest.py @@ -0,0 +1,26 @@ +#!/usr/bin/env python3 + +import json +import pytest +import pandas as pd + + +def pytest_addoption(parser): + parser.addoption("--input", action="store", help="Path to csv file.") + parser.addoption("--input-dir", action="store", help="Path to output dir.") + + +@pytest.fixture +def input_data(request): + filename = request.config.getoption("--input") + if filename: + with open(filename, "r") as inp: + return pd.read_csv(filename) + else: + return None + + +@pytest.fixture +def input_dir(request): + dirname = request.config.getoption("--input-dir") + return dirname diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/CMakeLists.txt new file mode 100644 index 0000000000..5122c328b1 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/CMakeLists.txt @@ -0,0 +1,51 @@ +# +# +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +project( + rocprofiler-tests-counter-collection + LANGUAGES CXX + VERSION 0.0.0) + +find_package(rocprofiler-sdk REQUIRED) + +foreach(FILENAME validate.py input.txt) + configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} + ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY) +endforeach() + +# pmc1 +add_test( + NAME rocprofv3-test-counter-collection-pmc1-execute + COMMAND + $ -i + ${CMAKE_CURRENT_BINARY_DIR}/input.txt -d ${CMAKE_CURRENT_BINARY_DIR}/out_cc_1 -o + pmc1 $) + +set(cc-env-pmc1 + "${PRELOAD_ENV}" + "HSA_TOOLS_LIB=$" + "LD_LIBRARY_PATH=$:$ENV{LD_LIBRARY_PATH}" + ) + +set_tests_properties( + rocprofv3-test-counter-collection-pmc1-execute + PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT "${cc-env-pmc1}" + FAIL_REGULAR_EXPRESSION "threw an exception") + +add_test(NAME rocprofv3-test-counter-collection-pmc1-validate + COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --input + "${CMAKE_CURRENT_BINARY_DIR}/out_cc_1/pmc_1/pmc1_counter_collection.csv") + +set_tests_properties( + rocprofv3-test-counter-collection-pmc1-validate + PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS + rocprofv3-test-counter-collection-pmc1-execute FAIL_REGULAR_EXPRESSION + "threw an exception") + +# Needs to be enabled once counter colelction tool support is merged +set_tests_properties(rocprofv3-test-counter-collection-pmc1-execute PROPERTIES DISABLED + TRUE) +set_tests_properties(rocprofv3-test-counter-collection-pmc1-validate PROPERTIES DISABLED + TRUE) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/input.txt b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/input.txt new file mode 100644 index 0000000000..4b4b2e8d41 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/input.txt @@ -0,0 +1 @@ +pmc: SQ_WAVES \ No newline at end of file diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/validate.py new file mode 100644 index 0000000000..ac4514d916 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input1/validate.py @@ -0,0 +1,16 @@ +import pandas as pd +import sys +import pytest + + +def test_validate_counter_collection_pmc1(input_data: pd.DataFrame): + df = input_data + + assert df.empty == False + assert df["agent-id"].map(type).eq(int).all() + assert len(df["kernel-name"]) > 0 + + +if __name__ == "__main__": + exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) + sys.exit(exit_code) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/CMakeLists.txt new file mode 100644 index 0000000000..fa8a1c812d --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/CMakeLists.txt @@ -0,0 +1,50 @@ +# +# +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +project( + rocprofiler-tests-counter-collection + LANGUAGES CXX + VERSION 0.0.0) + +find_package(rocprofiler-sdk REQUIRED) + +foreach(FILENAME validate.py input.txt) + configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} + ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY) +endforeach() + +# pmc2 +add_test( + NAME rocprofv3-test-counter-collection-pmc2-execute + COMMAND + $ -i + ${CMAKE_CURRENT_BINARY_DIR}/input.txt -d ${CMAKE_CURRENT_BINARY_DIR}/out_cc_2 -o + pmc2 $) +set(cc-env-pmc2 + "${PRELOAD_ENV}" + "HSA_TOOLS_LIB=$" + "LD_LIBRARY_PATH=$:$ENV{LD_LIBRARY_PATH}" + ) + +set_tests_properties( + rocprofv3-test-counter-collection-pmc2-execute + PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT "${cc-env-pmc2}" + FAIL_REGULAR_EXPRESSION "threw an exception") + +add_test(NAME rocprofv3-test-counter-collection-pmc2-validate + COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py + --input-dir "${CMAKE_CURRENT_BINARY_DIR}/out_cc_2") + +set_tests_properties( + rocprofv3-test-counter-collection-pmc2-validate + PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS + rocprofv3-test-counter-collection-pmc2-execute FAIL_REGULAR_EXPRESSION + "threw an exception") + +# Needs to be enabled once counter colelction tool support is merged +set_tests_properties(rocprofv3-test-counter-collection-pmc2-execute PROPERTIES DISABLED + TRUE) +set_tests_properties(rocprofv3-test-counter-collection-pmc2-validate PROPERTIES DISABLED + TRUE) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/input.txt b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/input.txt new file mode 100644 index 0000000000..a99a8d4730 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/input.txt @@ -0,0 +1,2 @@ +pmc: SQ_WAVES +pmc: GRBM_COUNT \ No newline at end of file diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/validate.py new file mode 100644 index 0000000000..4a9793d198 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/input2/validate.py @@ -0,0 +1,44 @@ +import pandas as pd +import os +import sys +import pytest + + +def test_validate_counter_collection_pmc2(input_dir: pd.DataFrame): + directory_path = input_dir + + # Check if the directory is not empty + assert os.path.isdir(directory_path), f"{directory_path} is not a directory." + assert os.listdir(directory_path), f"{directory_path} is empty." + + # Check if there are 2 subdirectories pmc_1 and pmc_2 + subdirectories = [ + d + for d in os.listdir(directory_path) + if os.path.isdir(os.path.join(directory_path, d)) + ] + assert ( + len(subdirectories) == 2 + ), f"Expected 2 subdirectories, found {len(subdirectories)}." + + # Check if each subdirectory has files + for subdirectory in subdirectories: + subdirectory_path = os.path.join(directory_path, subdirectory) + assert os.listdir(subdirectory_path), f"{subdirectory_path} is empty." + + # Check if each file in the subdirectory has some data + for file_name in os.listdir(subdirectory_path): + file_path = os.path.join(subdirectory_path, file_name) + assert os.path.isfile(file_path), f"{file_path} is not a file." + + with open(file_path, "r") as file: + df = pd.read_csv(file) + # check if kernel-name is present + assert len(df["kernel-name"]) > 0 + # check if counter value is positive + assert len(df["counter_value"]) > 0 + + +if __name__ == "__main__": + exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) + sys.exit(exit_code) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/pytest.ini b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/pytest.ini new file mode 100644 index 0000000000..75dce2e9c2 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/counter-collection/pytest.ini @@ -0,0 +1,5 @@ + +[pytest] +addopts = --durations=20 -ras -vv +testpaths = input1/validate.py + input2/validate.py diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt new file mode 100644 index 0000000000..fc6a6ec2eb --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt @@ -0,0 +1,49 @@ +# +# +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +project( + rocprofiler-tests-counter-collection + LANGUAGES CXX + VERSION 0.0.0) + +find_package(rocprofiler-sdk REQUIRED) + +foreach(FILENAME validate.py conftest.py pytest.ini input.txt) + configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} + ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY) +endforeach() + +# pmc3 +add_test( + NAME rocprofv3-test-tracing-plus-cc-execute + COMMAND + $ --hsa-trace -i + ${CMAKE_CURRENT_BINARY_DIR}/input.txt -d ${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace + -o pmc3 $) + +set(cc-tracing-env + "${PRELOAD_ENV}" + "HSA_TOOLS_LIB=$" + "LD_LIBRARY_PATH=$:$ENV{LD_LIBRARY_PATH}" + ) + +set_tests_properties( + rocprofv3-test-tracing-plus-cc-execute + PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT "${cc-tracing-env}" + FAIL_REGULAR_EXPRESSION "threw an exception") + +add_test(NAME rocprofv3-test-tracing-plus-cc-validate + COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py + --input-dir "${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace") + +set_tests_properties( + rocprofv3-test-tracing-plus-cc-validate + PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS + rocprofv3-test-tracing-plus-cc-execute FAIL_REGULAR_EXPRESSION + "threw an exception") + +# Needs to be enabled once counter colelction tool support is merged +set_tests_properties(rocprofv3-test-tracing-plus-cc-execute PROPERTIES DISABLED TRUE) +set_tests_properties(rocprofv3-test-tracing-plus-cc-validate PROPERTIES DISABLED TRUE) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/conftest.py new file mode 100644 index 0000000000..ad013dbe3a --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/conftest.py @@ -0,0 +1,15 @@ +#!/usr/bin/env python3 + +import json +import pytest +import pandas as pd + + +def pytest_addoption(parser): + parser.addoption("--input-dir", action="store", help="Path to output dir.") + + +@pytest.fixture +def input_dir(request): + dirname = request.config.getoption("--input-dir") + return dirname diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/input.txt b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/input.txt new file mode 100644 index 0000000000..2c1d6151b6 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/input.txt @@ -0,0 +1,6 @@ +# multi block counters + +pmc: GRBM_COUNT GRBM_GUI_ACTIVE SQ_CYCLES SQ_BUSY_CYCLES SQ_WAVES +pmc: TCC_REQ_sum TCC_STREAMING_REQ_sum TCC_HIT_sum TCC_MISS_sum +pmc: TCC_EA_WRREQ_sum TCC_EA_WRREQ_64B_sum TCC_EA_WR_UNCACHED_32B_sum +pmc: TCC_TOO_MANY_EA_WRREQS_STALL_sum TCC_EA_ATOMIC_sum TCC_EA_RDREQ_sum TCC_EA_RDREQ_32B_sum \ No newline at end of file diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/pytest.ini b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/pytest.ini new file mode 100644 index 0000000000..899bc38408 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/pytest.ini @@ -0,0 +1,4 @@ + +[pytest] +addopts = --durations=20 -ras -vv +testpaths = validate.py diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/validate.py new file mode 100644 index 0000000000..ef54f9f781 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-plus-cc/validate.py @@ -0,0 +1,42 @@ +import pandas as pd +import os +import sys +import pytest + + +def test_validate_counter_collection_plus_tracing(input_dir: pd.DataFrame): + directory_path = input_dir + + # Check if the directory is not empty + assert os.path.isdir(directory_path), f"{directory_path} is not a directory." + assert os.listdir(directory_path), f"{directory_path} is empty." + + # Check if there are 4 subdirectories for pmc's + subdirectories = [ + d + for d in os.listdir(directory_path) + if os.path.isdir(os.path.join(directory_path, d)) + ] + assert ( + len(subdirectories) == 4 + ), f"Expected 4 subdirectories, found {len(subdirectories)}." + + # Check if each subdirectory has files + for subdirectory in subdirectories: + subdirectory_path = os.path.join(directory_path, subdirectory) + assert os.listdir(subdirectory_path), f"{subdirectory_path} is empty." + + # Check if each file in the subdirectory has some data + for file_name in os.listdir(subdirectory_path): + file_path = os.path.join(subdirectory_path, file_name) + assert os.path.isfile(file_path), f"{file_path} is not a file." + + with open(file_path, "r") as file: + df = pd.read_csv(file) + # check if either kernel-name/FUNCTION is present + assert "kernel-name" in df.columns or "FUNCTION" in df.columns + + +if __name__ == "__main__": + exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) + sys.exit(exit_code) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/CMakeLists.txt new file mode 100644 index 0000000000..57dd050ef3 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/CMakeLists.txt @@ -0,0 +1,47 @@ +# +# rocprofv3 tool test +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +project( + rocprofiler-tests-rocprofv3-tracing + LANGUAGES CXX + VERSION 0.0.0) + +find_package(rocprofiler-sdk REQUIRED) + +add_test( + NAME rocprofv3-test-trace-execute + COMMAND + $ --hsa-trace --kernel-trace + --memory-copy-trace -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out + $) + +string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV + "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") + +set(tracing-env "${PRELOAD_ENV}" + "HSA_TOOLS_LIB=$") + +set_tests_properties(rocprofv3-test-trace-execute PROPERTIES LABELS "integration-tests" + ENVIRONMENT "${tracing-env}") + +foreach(FILENAME validate.py conftest.py) + configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} + ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY) +endforeach() + +add_test( + NAME rocprofv3-test-trace-validate + COMMAND + ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --hsa-input + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_hsa_api_trace.csv + --kernel-input + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_kernel_trace.csv + --memory-copy-input + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_memory_copy_trace.csv) + +set_tests_properties( + rocprofv3-test-trace-validate + PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS rocprofv3-test-trace-execute + FAIL_REGULAR_EXPRESSION "AssertionError") diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/conftest.py new file mode 100644 index 0000000000..febffdf809 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/conftest.py @@ -0,0 +1,58 @@ +#!/usr/bin/env python3 + +import csv +import pytest + + +def pytest_addoption(parser): + parser.addoption( + "--hsa-input", + action="store", + help="Path to HSA API tracing CSV file.", + ) + parser.addoption( + "--kernel-input", + action="store", + help="Path to kernel tracing CSV file.", + ) + parser.addoption( + "--memory-copy-input", + action="store", + help="Path to memory-copy tracing CSV file.", + ) + + +@pytest.fixture +def hsa_input_data(request): + filename = request.config.getoption("--hsa-input") + data = [] + with open(filename, "r") as inp: + reader = csv.DictReader(inp) + for row in reader: + data.append(row) + + return data + + +@pytest.fixture +def kernel_input_data(request): + filename = request.config.getoption("--kernel-input") + data = [] + with open(filename, "r") as inp: + reader = csv.DictReader(inp) + for row in reader: + data.append(row) + + return data + + +@pytest.fixture +def memory_copy_input_data(request): + filename = request.config.getoption("--memory-copy-input") + data = [] + with open(filename, "r") as inp: + reader = csv.DictReader(inp) + for row in reader: + data.append(row) + + return data diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py new file mode 100644 index 0000000000..e051a39fb3 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py @@ -0,0 +1,68 @@ +#!/usr/bin/env python3 + +import sys +import pytest + + +def test_hsa_api_trace(hsa_input_data): + functions = [] + correlation_ids = [] + for row in hsa_input_data: + assert row["Domain"] == "HSA_API" + assert int(row["Process_Id"]) > 0 + assert int(row["Thread_Id"]) >= int(row["Process_Id"]) + assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"]) + functions.append(row["Function"]) + correlation_ids.append(int(row["Correlation_Id"])) + + correlation_ids = sorted(list(set(correlation_ids))) + + # all correlation ids are unique + assert len(correlation_ids) == len(hsa_input_data) + # correlation ids are numbered from 1 to N + assert correlation_ids[0] == 1 + assert correlation_ids[-1] == len(correlation_ids) + + functions = list(set(functions)) + assert "hsa_amd_memory_async_copy_on_engine" in functions + + +def test_kernel_trace(kernel_input_data): + for row in kernel_input_data: + assert row["Kind"] == "KERNEL_DISPATCH" + assert int(row["Agent_Id"]) > 0 + assert int(row["Queue_Id"]) > 0 + assert int(row["Kernel_Id"]) > 0 + assert row["Kernel_Name"] == "matrixTranspose(float*, float*, int)" + assert int(row["Correlation_Id"]) > 0 + assert int(row["Workgroup_Size_X"]) == 4 + assert int(row["Workgroup_Size_Y"]) == 4 + assert int(row["Workgroup_Size_Z"]) == 1 + assert int(row["Grid_Size_X"]) == 1024 + assert int(row["Grid_Size_Y"]) == 1024 + assert int(row["Grid_Size_Z"]) == 1 + assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"]) + + +def test_memory_copy_trace(memory_copy_input_data): + for row in memory_copy_input_data: + assert row["Kind"] == "MEMORY_COPY" + + row = memory_copy_input_data[0] + assert row["Direction"] == "HOST_TO_DEVICE" + assert int(row["Source_Agent_Id"]) == 0 + assert int(row["Destination_Agent_Id"]) == 1 + assert int(row["Correlation_Id"]) > 0 + assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"]) + + row = memory_copy_input_data[1] + assert row["Direction"] == "DEVICE_TO_HOST" + assert int(row["Source_Agent_Id"]) == 1 + assert int(row["Destination_Agent_Id"]) == 0 + assert int(row["Correlation_Id"]) > 0 + assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"]) + + +if __name__ == "__main__": + exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) + sys.exit(exit_code) diff --git a/projects/rocprofiler-sdk/tests/tools/json-tool.cpp b/projects/rocprofiler-sdk/tests/tools/json-tool.cpp index ce93f60a08..0ceebb34ae 100644 --- a/projects/rocprofiler-sdk/tests/tools/json-tool.cpp +++ b/projects/rocprofiler-sdk/tests/tools/json-tool.cpp @@ -181,7 +181,9 @@ buffer_name_info get_buffer_tracing_names() { static const auto supported = std::unordered_set{ - ROCPROFILER_BUFFER_TRACING_HSA_API, ROCPROFILER_BUFFER_TRACING_MARKER_API}; + ROCPROFILER_BUFFER_TRACING_HSA_API, + ROCPROFILER_BUFFER_TRACING_MARKER_API, + ROCPROFILER_BUFFER_TRACING_MEMORY_COPY}; auto cb_name_info = buffer_name_info{}; //