diff --git a/CHANGELOG.md b/CHANGELOG.md index f0552d4417..8be19572b9 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -244,6 +244,20 @@ Rocprofiler for ROCm 5.7 added support for counter collection (PMC) and advanced - "--mode file" option in ATT, which allows for parsed files to be stored. Run python3 httpserver.py from within ./UI/ to view files locally. - "ROCPROFILER_MAX_ATT_PROFILES" environment variable can be set. Previously fixed at 16, now the default is 1. - Increased ATT buffer size per collection to 1GB. +- File plugin is splitted to File & CLI plugins, CLI plugin is responsible for showing results on the terminal screen and will be automatically the choice if no -d option given in rocprof, File plugin on the other hand is responsible for writing the output results in files if -d option is given. +- Structure of the results is different for both CLI & File plugin; File plugin will make sure every type of result is in a separate file, starting by specifying the header; CLI plugin will have header for kernel dispatches and counter collection and another header for tracing results; Example: + ``` + Dispatch_ID,GPU_ID,Queue_ID,Queue_Index,PID,TID,GRD,WGR,LDS,SCR,Arch_VGPR,ACCUM_VGPR,SGPR,Wave_Size,SIG,OBJ,Kernel_Name,Start_Timestamp,End_Timestamp,Counters + + 1,4,1,1,1584730,1584730,10,10,0,0,8,0,16,64,140464978048000,1,"helloworld(char*, char*) (.kd)",0,140469300947216,GRBM_COUNT,12637.000000 + ``` + ``` + Record_ID,Domain,Function,Operation,Kernel_Name,Start_Timestamp,End_Timestamp,Correlation_ID,ROCTX_ID,ROCTX_Message + + 2,HIP_API_DOMAIN,hipGetDeviceProperties,,,316678074094190,316678074098929,1,, + 4,HIP_API_DOMAIN,hipMalloc,,,316678074105702,316678074130851,2,, + 6,HIP_API_DOMAIN,hipMalloc,,,316678074131382,316678074136111,3,, + ``` ### Fixed - Samples are fixed to show the new usage of phases. - Plugin option validates the plugin names. diff --git a/bin/rocprofv2 b/bin/rocprofv2 index 4208e5b3ca..f4f5fdac05 100755 --- a/bin/rocprofv2 +++ b/bin/rocprofv2 @@ -267,6 +267,9 @@ COUNTERS_PMC_DIRS="" if [ -n "$PMC_LINES" ]; then COUNTER=1 for i in ${!PMC_LINES[@]}; do + if [[ "${PMC_LINES[$i]}" != *"pmc:"* ]]; then + continue + fi export ROCPROFILER_COUNTERS="${PMC_LINES[$i]}" if [ -n "$OUTPUT_PATH" ]; then if [ ! -n "$ATT_ARGV" ]; then diff --git a/include/rocprofiler/v2/rocprofiler_plugin.h b/include/rocprofiler/v2/rocprofiler_plugin.h index 882019c48b..516e2f9b17 100644 --- a/include/rocprofiler/v2/rocprofiler_plugin.h +++ b/include/rocprofiler/v2/rocprofiler_plugin.h @@ -77,10 +77,12 @@ extern "C" { * \p rocprofiler_major_version matches and this is greater than the minor * version of the ROCProfiler API used to build the plugin library. This ensures * compatibility of the trace data format. + * @param[in] data Pointer to the data passed to the ROCProfiler Plugin by the tool * @return Returns 0 on success and -1 on error. */ ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_version, - uint32_t rocprofiler_minor_version); + uint32_t rocprofiler_minor_version, + void* data); /** * Finalize plugin. diff --git a/plugin/CMakeLists.txt b/plugin/CMakeLists.txt index 1576b72a9b..922a8b1862 100644 --- a/plugin/CMakeLists.txt +++ b/plugin/CMakeLists.txt @@ -24,3 +24,4 @@ add_subdirectory(file) add_subdirectory(perfetto) add_subdirectory(ctf) add_subdirectory(att) +add_subdirectory(cli) diff --git a/plugin/att/att.cpp b/plugin/att/att.cpp index 046bdd100a..e9995dee42 100644 --- a/plugin/att/att.cpp +++ b/plugin/att/att.cpp @@ -79,9 +79,10 @@ class att_plugin_t { CHECK_ROCPROFILER(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME, att_tracer_record->kernel_id, &kernel_name_c)); - std::string name_demangled = rocprofiler::truncate_name(rocprofiler::cxx_demangle(kernel_name_c)); + std::string name_demangled = + rocprofiler::truncate_name(rocprofiler::cxx_demangle(kernel_name_c)); - if (name_demangled.size() > ATT_FILENAME_MAXBYTES) { // Limit filename size + if (name_demangled.size() > ATT_FILENAME_MAXBYTES) { // Limit filename size name_demangled = name_demangled.substr(0, ATT_FILENAME_MAXBYTES); } @@ -167,7 +168,8 @@ att_plugin_t* att_plugin = nullptr; } // namespace ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_version, - uint32_t rocprofiler_minor_version) { + uint32_t rocprofiler_minor_version, + void* data) { if (rocprofiler_major_version != ROCPROFILER_VERSION_MAJOR || rocprofiler_minor_version < ROCPROFILER_VERSION_MINOR) return -1; diff --git a/plugin/cli/CMakeLists.txt b/plugin/cli/CMakeLists.txt new file mode 100644 index 0000000000..6a4098375d --- /dev/null +++ b/plugin/cli/CMakeLists.txt @@ -0,0 +1,47 @@ +# ############################################################################### +# # Copyright (c) 2022 Advanced Micro Devices, Inc. +# # +# # Permission is hereby granted, free of charge, to any person obtaining a copy +# # of this software and associated documentation files (the "Software"), to +# # deal in the Software without restriction, including without limitation the +# # rights to use, copy, modify, merge, publish, distribute, sublicense, and/or +# # sell copies of the Software, and to permit persons to whom the Software is +# # furnished to do so, subject to the following conditions: +# # +# # The above copyright notice and this permission notice shall be included in +# # all copies or substantial portions of the Software. +# # +# # THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# # IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# # FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# # AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# # LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING +# # FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS +# # IN THE SOFTWARE. +# ############################################################################### + +file(GLOB ROCPROFILER_UTIL_SRC_FILES ${PROJECT_SOURCE_DIR}/src/utils/helper.cpp) + +file(GLOB CLI_SOURCES "*.cpp") +add_library(cli_plugin SHARED ${CLI_SOURCES} ${ROCPROFILER_UTIL_SRC_FILES}) + +set_target_properties(cli_plugin PROPERTIES + CXX_VISIBILITY_PRESET hidden + LINK_DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/../exportmap + LIBRARY_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR}) + +target_compile_definitions(cli_plugin + PRIVATE HIP_PROF_HIP_API_STRING=1 __HIP_PLATFORM_HCC__=1) + +target_include_directories(cli_plugin PRIVATE ${PROJECT_SOURCE_DIR}) + +target_link_options(cli_plugin PRIVATE -Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/../exportmap -Wl,--no-undefined) + +target_link_libraries(cli_plugin PRIVATE rocprofiler-v2 hsa-runtime64::hsa-runtime64 stdc++fs atomic amd_comgr dl) + +install(TARGETS cli_plugin LIBRARY + DESTINATION ${CMAKE_INSTALL_LIBDIR}/${PROJECT_NAME} + COMPONENT asan) +install(TARGETS cli_plugin LIBRARY + DESTINATION ${CMAKE_INSTALL_LIBDIR}/${PROJECT_NAME} + COMPONENT runtime) diff --git a/plugin/cli/cli.cpp b/plugin/cli/cli.cpp new file mode 100644 index 0000000000..8a336892cd --- /dev/null +++ b/plugin/cli/cli.cpp @@ -0,0 +1,484 @@ +/* Copyright (c) 2022 Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. */ + +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "rocprofiler.h" +#include "rocprofiler_plugin.h" +#include "../utils.h" + +namespace fs = std::experimental::filesystem; + +namespace { + +static std::string output_file_name; +class file_plugin_t { + private: + enum class output_type_t { COUNTER, TRACER, PC_SAMPLING }; + + class output_file_t { + public: + output_file_t(std::string name) : name_(std::move(name)) {} + + std::string name() const { return name_; } + + template std::ostream& operator<<(T&& value) { + if (!is_open()) open(); + return stream_ << std::forward(value); + } + + std::ostream& operator<<(std::ostream& (*func)(std::ostream&)) { + if (!is_open()) open(); + return stream_ << func; + } + + void open() { + // If the stream is already in the failed state, there's no need to try + // to open the file. + if (fail()) return; + + const char* output_dir = getenv("OUTPUT_PATH"); + output_file_name = getenv("OUT_FILE_NAME") ? std::string(getenv("OUT_FILE_NAME")) + "_" : ""; + + if (output_dir == nullptr && getenv("OUT_FILE_NAME") == nullptr) { + stream_.copyfmt(std::cout); + stream_.clear(std::cout.rdstate()); + stream_.basic_ios::rdbuf(std::cout.rdbuf()); + return; + } + if (output_dir == nullptr) output_dir = "./"; + + fs::path output_prefix(output_dir); + if (!fs::is_directory(fs::status(output_prefix))) { + if (!stream_.fail()) rocprofiler::warning("Cannot open output directory '%s'", output_dir); + stream_.setstate(std::ios_base::failbit); + return; + } + + std::stringstream ss; + output_file_name = replace_MPI_macros(output_file_name); + + ss << output_file_name << GetPid() << "_" << name_; + stream_.open(output_prefix / ss.str()); + } + + bool is_open() const { return stream_.is_open(); } + bool fail() const { return stream_.fail(); } + + // Returns a string with the MPI %macro replaced with the corresponding envvar + std::string replace_MPI_macros(std::string output_file_name) { + std::unordered_map MPI_BUILTINS = { + {"MPI_RANK", "%rank"}, + {"OMPI_COMM_WORLD_RANK", "%rank"}, + {"MV2_COMM_WORLD_RANK", "%rank"}}; + + for (const auto& [envvar, key] : MPI_BUILTINS) { + size_t key_find = output_file_name.rfind(key); + if (key_find == std::string::npos) continue; // Does not contain a %?rank var + + const char* env_var_set = getenv(envvar); + if (env_var_set == nullptr) continue; // MPI_COMM_WORLD_x var is does not exist + + int rank = atoi(env_var_set); + output_file_name = output_file_name.substr(0, key_find) + std::to_string(rank) + + output_file_name.substr(key_find + std::string(key).size()); + } + + return output_file_name; + } + + private: + const std::string name_; + std::ofstream stream_; + }; + + output_file_t* get_output_file(output_type_t output_type, uint32_t domain = 0) { + switch (output_type) { + case output_type_t::COUNTER: + return &output_file_; + case output_type_t::TRACER: + switch (domain) { + case ACTIVITY_DOMAIN_ROCTX: + return &roctx_file_; + case ACTIVITY_DOMAIN_HSA_API: + return &hsa_api_file_; + case ACTIVITY_DOMAIN_HIP_API: + return &hip_api_file_; + case ACTIVITY_DOMAIN_HIP_OPS: + return &hip_activity_file_; + case ACTIVITY_DOMAIN_HSA_OPS: + return &hsa_async_copy_file_; + default: + assert(!"domain/op not supported!"); + break; + } + break; + case output_type_t::PC_SAMPLING: + return &pc_sample_file_; + } + return nullptr; + } + + public: + file_plugin_t() { + valid_ = true; + } + + std::mutex writing_lock; + + const char* GetDomainName(rocprofiler_tracer_activity_domain_t domain) { + switch (domain) { + case ACTIVITY_DOMAIN_ROCTX: + return "ROCTX_DOMAIN"; + break; + case ACTIVITY_DOMAIN_HIP_API: + return "HIP_API_DOMAIN"; + break; + case ACTIVITY_DOMAIN_HIP_OPS: + return "HIP_OPS_DOMAIN"; + break; + case ACTIVITY_DOMAIN_HSA_API: + return "HSA_API_DOMAIN"; + break; + case ACTIVITY_DOMAIN_HSA_OPS: + return "HSA_OPS_DOMAIN"; + break; + case ACTIVITY_DOMAIN_HSA_EVT: + return "HSA_EVT_DOMAIN"; + break; + default: + return ""; + } + } + + void FlushTracerRecord(rocprofiler_record_tracer_t tracer_record, + rocprofiler_session_id_t session_id, + rocprofiler_buffer_id_t buffer_id = rocprofiler_buffer_id_t{0}) { + std::lock_guard lock(writing_lock); + if (tracer_record.timestamps.end.value <= 0 && tracer_record.domain != ACTIVITY_DOMAIN_ROCTX) + return; + std::string function_name; + std::string kernel_name; + std::string roctx_message; + uint64_t roctx_id; + if ((tracer_record.operation_id.id == 0 && tracer_record.domain == ACTIVITY_DOMAIN_HIP_OPS)) { + if (tracer_record.name) { + kernel_name = rocprofiler::cxx_demangle(tracer_record.name); + std::string key = "\""; + std::size_t found = kernel_name.rfind(key); + while (found != std::string::npos) { + kernel_name.replace(found, key.length(), "'"); + found = kernel_name.rfind(key, found - 1); + } + } + } + size_t function_name_size = 0; + char* function_name_c = nullptr; + if (tracer_record.domain == ACTIVITY_DOMAIN_HSA_API) { + CHECK_ROCPROFILER(rocprofiler_query_hsa_tracer_api_data_info_size( + rocprofiler_session_id_t{0}, ROCPROFILER_HSA_FUNCTION_NAME, tracer_record.api_data_handle, + tracer_record.operation_id, &function_name_size)); + function_name_c = new char[function_name_size]; + if (function_name_size > 1) { + CHECK_ROCPROFILER(rocprofiler_query_hsa_tracer_api_data_info( + rocprofiler_session_id_t{0}, ROCPROFILER_HSA_FUNCTION_NAME, + tracer_record.api_data_handle, tracer_record.operation_id, &function_name_c)); + if (function_name_c) function_name = std::string(function_name_c); + } + } + if (tracer_record.domain == ACTIVITY_DOMAIN_HIP_API) { + CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info_size( + rocprofiler_session_id_t{0}, ROCPROFILER_HIP_FUNCTION_NAME, tracer_record.api_data_handle, + tracer_record.operation_id, &function_name_size)); + if (function_name_size > 1) { + CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info( + session_id, ROCPROFILER_HIP_FUNCTION_NAME, tracer_record.api_data_handle, + tracer_record.operation_id, &function_name_c)); + if (function_name_c) function_name = std::string(function_name_c); + } + if (tracer_record.name) { + kernel_name = rocprofiler::cxx_demangle(std::string(tracer_record.name)); + std::string key = "\""; + std::size_t found = kernel_name.rfind(key); + while (found != std::string::npos) { + kernel_name.replace(found, key.length(), "'"); + found = kernel_name.rfind(key, found - 1); + } + // TODO: Change how this API returns a string. + } + } + if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX) { + if (tracer_record.name) roctx_message = rocprofiler::cxx_demangle(tracer_record.name); + roctx_id = tracer_record.operation_id.id; + } + char* activity_name = nullptr; + if (tracer_record.domain == ACTIVITY_DOMAIN_HIP_OPS) { + if (tracer_record.api_data_handle.handle) { + kernel_name = rocprofiler::cxx_demangle( + const_cast(reinterpret_cast(tracer_record.api_data_handle.handle))); + } + size_t activity_name_size = 0; + CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info_size( + session_id, ROCPROFILER_HIP_ACTIVITY_NAME, tracer_record.api_data_handle, + tracer_record.operation_id, &activity_name_size)); + if (activity_name_size > 1) { + activity_name = nullptr; + CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info( + session_id, ROCPROFILER_HIP_ACTIVITY_NAME, tracer_record.api_data_handle, + tracer_record.operation_id, &activity_name)); + } + } + if (tracer_record.domain == ACTIVITY_DOMAIN_HSA_OPS) { + size_t activity_name_size = 0; + CHECK_ROCPROFILER(rocprofiler_query_hsa_tracer_api_data_info_size( + session_id, ROCPROFILER_HSA_ACTIVITY_NAME, tracer_record.api_data_handle, + tracer_record.operation_id, &activity_name_size)); + if (activity_name_size > 1) { + activity_name = nullptr; + CHECK_ROCPROFILER(rocprofiler_query_hsa_tracer_api_data_info( + session_id, ROCPROFILER_HSA_ACTIVITY_NAME, tracer_record.api_data_handle, + tracer_record.operation_id, &activity_name)); + } + } + + output_file_t* output_file = get_output_file(output_type_t::TRACER, tracer_record.domain); + *output_file << "Record_ID(" << tracer_record.header.id.handle << "), " + << "Domain(" << GetDomainName(tracer_record.domain) << "), "; + if (function_name.size() > 1) *output_file << "Function(" << function_name << "), "; + if (activity_name) *output_file << "Operation_Name(" << activity_name << "), "; + if (kernel_name.size() > 1) *output_file << "Kernel_Name(" << kernel_name.c_str() << "), "; + if (tracer_record.domain != ACTIVITY_DOMAIN_ROCTX) { + *output_file << "Start_Timestamp(" << tracer_record.timestamps.begin.value << "), " + << "End_Timestamp(" << tracer_record.timestamps.end.value << "), " + << "Correlation_ID(" << tracer_record.correlation_id.value << ")"; + } else { + *output_file << "Timestamp(" << tracer_record.timestamps.begin.value << "), "; + } + if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX && roctx_id >= 0) + *output_file << "ROCTX_ID(" << roctx_id << "), "; + if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX && roctx_message.size() > 1) + *output_file << "ROCTX_Message(" << roctx_message << ")"; + *output_file << std::endl; + } + + void FlushProfilerRecord(const rocprofiler_record_profiler_t* profiler_record, + rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) { + std::lock_guard lock(writing_lock); + size_t name_length = 0; + output_file_t* output_file{nullptr}; + output_file = get_output_file(output_type_t::COUNTER); + CHECK_ROCPROFILER(rocprofiler_query_kernel_info_size(ROCPROFILER_KERNEL_NAME, + profiler_record->kernel_id, &name_length)); + // Taken from rocprofiler: The size hasn't changed in recent past + static const uint32_t lds_block_size = 128 * 4; + const char* kernel_name_c = nullptr; + if (name_length > 1) { + CHECK_ROCPROFILER(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME, + profiler_record->kernel_id, &kernel_name_c)); + } + *output_file << "Record_ID(" << std::to_string(profiler_record->header.id.handle) << "), " + << "GPU_ID(" << std::to_string(profiler_record->gpu_id.handle) << "), " + << "Queue_ID(" << std::to_string(profiler_record->queue_id.handle) << "), " + << "Queue_Index(" << std::to_string(profiler_record->queue_idx.value) << "), " + << "Process_ID(" << std::to_string(GetPid()) << "), " + << "Thread_ID(" << std::to_string(profiler_record->thread_id.value) << "), " + << "Grid_Size(" << std::to_string(profiler_record->kernel_properties.grid_size) + << "), " + << "Workgroup_Size(" + << std::to_string(profiler_record->kernel_properties.workgroup_size) << "), " + << "LDS(" + << std::to_string( + ((profiler_record->kernel_properties.lds_size + (lds_block_size - 1)) & + ~(lds_block_size - 1))) + << "), " + << "Scratch_Size(" + << std::to_string(profiler_record->kernel_properties.scratch_size) << "), " + << "Arch_VGPR(" + << std::to_string(profiler_record->kernel_properties.arch_vgpr_count) << "), " + << "Accumulative_VGPR(" + << std::to_string(profiler_record->kernel_properties.accum_vgpr_count) << "), " + << "SGPR(" << std::to_string(profiler_record->kernel_properties.sgpr_count) + << "), " + << "Wave_Size(" << std::to_string(profiler_record->kernel_properties.wave_size); + std::string kernel_name = ""; + if (name_length > 1) { + kernel_name = rocprofiler::truncate_name(rocprofiler::cxx_demangle(kernel_name_c)); + std::string key = "\""; + std::size_t found = kernel_name.rfind(key); + while (found != std::string::npos) { + kernel_name.replace(found, key.length(), "'"); + found = kernel_name.rfind(key, found - 1); + } + } + *output_file << "), " + << "Kernel_Name(\"" << kernel_name << "\"), " + << "Begin_Timestamp(" << std::to_string(profiler_record->timestamps.begin.value) + << "), " + << "End_Timestamp(" << std::to_string(profiler_record->timestamps.end.value) + << ")"; + + // For Counters + if (profiler_record->counters) { + *output_file << ", "; + for (uint64_t i = 0; i < profiler_record->counters_count.value; i++) { + if (profiler_record->counters[i].counter_handler.handle > 0) { + size_t counter_name_length = 0; + CHECK_ROCPROFILER(rocprofiler_query_counter_info_size( + session_id, ROCPROFILER_COUNTER_NAME, profiler_record->counters[i].counter_handler, + &counter_name_length)); + if (counter_name_length > 1) { + const char* name_c = nullptr; + CHECK_ROCPROFILER(rocprofiler_query_counter_info( + session_id, ROCPROFILER_COUNTER_NAME, profiler_record->counters[i].counter_handler, + &name_c)); + *output_file << name_c << "(" + << std::to_string(profiler_record->counters[i].value.value) << ")"; + if (i < profiler_record->counters_count.value - 1) *output_file << ", "; + } + } + } + } + *output_file << std::endl; + if (kernel_name_c) { + free(const_cast(kernel_name_c)); + } + } + + void FlushPCSamplingRecord(const rocprofiler_record_pc_sample_t* pc_sampling_record) { + output_file_t* output_file{nullptr}; + output_file = get_output_file(output_type_t::PC_SAMPLING); + const auto& sample = pc_sampling_record->pc_sample; + *output_file << "dispatch[" << sample.dispatch_id.value << "], " + << "timestamp(" << sample.timestamp.value << "), " + << "gpu_id(" << sample.gpu_id.handle << "), " + << "pc-sample(" << std::hex << std::showbase << sample.pc << "), " + << "se(" << sample.se << ')' << std::endl; + } + int WriteBufferRecords(const rocprofiler_record_header_t* begin, + const rocprofiler_record_header_t* end, + rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) { + while (begin < end) { + if (!begin) return 0; + switch (begin->kind) { + case ROCPROFILER_PROFILER_RECORD: { + const rocprofiler_record_profiler_t* profiler_record = + reinterpret_cast(begin); + FlushProfilerRecord(profiler_record, session_id, buffer_id); + break; + } + case ROCPROFILER_TRACER_RECORD: { + rocprofiler_record_tracer_t* tracer_record = const_cast( + reinterpret_cast(begin)); + FlushTracerRecord(*tracer_record, session_id, buffer_id); + break; + } + case ROCPROFILER_ATT_TRACER_RECORD: { + break; + } + case ROCPROFILER_PC_SAMPLING_RECORD: { + const rocprofiler_record_pc_sample_t* pc_sampling_record = + reinterpret_cast(begin); + FlushPCSamplingRecord(pc_sampling_record); + break; + } + default: + break; + } + rocprofiler_next_record(begin, &begin, session_id, buffer_id); + } + return 0; + } + + bool is_valid() const { return valid_; } + + private: + bool valid_{false}; + std::atomic tracer_header_written_{false}; + std::atomic profiler_header_written_{false}; + + output_file_t roctx_file_{"roctx_trace.txt"}, hsa_api_file_{"hsa_api_trace.txt"}, + hip_api_file_{"hip_api_trace.txt"}, hip_activity_file_{"hcc_ops_trace.txt"}, + hsa_async_copy_file_{"async_copy_trace.txt"}, pc_sample_file_{"pcs_trace.txt"}, + output_file_{"results.txt"}; +}; + +file_plugin_t* file_plugin = nullptr; + +} // namespace + +ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_version, + uint32_t rocprofiler_minor_version, + void* data) { + if (rocprofiler_major_version != ROCPROFILER_VERSION_MAJOR || + rocprofiler_minor_version < ROCPROFILER_VERSION_MINOR) + return -1; + + if (file_plugin != nullptr) return -1; + + file_plugin = new file_plugin_t(); + if (file_plugin->is_valid()) return 0; + + // The plugin failed to initialized, destroy it and return an error. + delete file_plugin; + file_plugin = nullptr; + return -1; +} + +ROCPROFILER_EXPORT void rocprofiler_plugin_finalize() { + if (!file_plugin) return; + delete file_plugin; + file_plugin = nullptr; +} + +ROCPROFILER_EXPORT int rocprofiler_plugin_write_buffer_records( + const rocprofiler_record_header_t* begin, const rocprofiler_record_header_t* end, + rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) { + if (!file_plugin || !file_plugin->is_valid()) return -1; + return file_plugin->WriteBufferRecords(begin, end, session_id, buffer_id); +} + +ROCPROFILER_EXPORT int rocprofiler_plugin_write_record(rocprofiler_record_tracer_t record) { + if (!file_plugin || !file_plugin->is_valid()) return -1; + if (record.header.id.handle == 0) return 0; + file_plugin->FlushTracerRecord(record, rocprofiler_session_id_t{0}, rocprofiler_buffer_id_t{0}); + return 0; +} diff --git a/plugin/ctf/ctf.cpp b/plugin/ctf/ctf.cpp index 13b98a36d6..5a90296ee5 100644 --- a/plugin/ctf/ctf.cpp +++ b/plugin/ctf/ctf.cpp @@ -38,7 +38,8 @@ rocm_ctf::Plugin* the_plugin = nullptr; } // namespace ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(const uint32_t rocprofiler_major_version, - const uint32_t rocprofiler_minor_version) { + const uint32_t rocprofiler_minor_version, + void* data) { if (rocprofiler_major_version != ROCPROFILER_VERSION_MAJOR || rocprofiler_minor_version < ROCPROFILER_VERSION_MINOR) { return -1; diff --git a/plugin/file/file.cpp b/plugin/file/file.cpp index 73135a3916..72125647af 100644 --- a/plugin/file/file.cpp +++ b/plugin/file/file.cpp @@ -26,6 +26,7 @@ #include #include +#include #include #include #include @@ -49,6 +50,36 @@ namespace fs = std::experimental::filesystem; namespace { +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; +} + static std::string output_file_name; class file_plugin_t { private: @@ -78,7 +109,7 @@ class file_plugin_t { if (fail()) return; const char* output_dir = getenv("OUTPUT_PATH"); - output_file_name = getenv("OUT_FILE_NAME") ? std::string(getenv("OUT_FILE_NAME")) + "_" : ""; + output_file_name = getenv("OUT_FILE_NAME") ? std::string(getenv("OUT_FILE_NAME")) : ""; if (output_dir == nullptr && getenv("OUT_FILE_NAME") == nullptr) { stream_.copyfmt(std::cout); @@ -96,10 +127,12 @@ class file_plugin_t { return; } - std::stringstream ss; output_file_name = replace_MPI_macros(output_file_name); - ss << output_file_name << GetPid() << "_" << name_; + std::stringstream ss; + ss << name_ << "_" << ((output_file_name.empty()) ? std::to_string(GetPid()) : "") + << output_file_name << ".csv"; + std::cout << "Results File: " << output_prefix / ss.str() << std::endl; stream_.open(output_prefix / ss.str()); } @@ -163,45 +196,90 @@ class file_plugin_t { } public: - file_plugin_t() { - output_file_t hsa_handles("hsa_handles.txt", true); - - [[maybe_unused]] hsa_status_t status = hsa_iterate_agents( - [](hsa_agent_t agent, void* user_data) { - auto* file = static_cast(user_data); - hsa_device_type_t type; - - if (hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type) != HSA_STATUS_SUCCESS) - return HSA_STATUS_ERROR; - - if (!file->isStdOut()) - *file << std::hex << std::showbase << agent.handle << " agent " - << ((type == HSA_DEVICE_TYPE_CPU) ? "cpu" : "gpu") << std::endl; - return HSA_STATUS_SUCCESS; - }, - &hsa_handles); - assert(status == HSA_STATUS_SUCCESS && "failed to iterate HSA agents"); - if (hsa_handles.fail()) { - rocprofiler::warning("Cannot write to '%s'", hsa_handles.name().c_str()); - return; - } - - // App begin timestamp begin_ts_file.txt - output_file_t begin_ts("begin_ts_file.txt", true); - - [[maybe_unused]] rocprofiler_timestamp_t app_begin_timestamp = {}; - CHECK_ROCPROFILER(rocprofiler_get_timestamp(&app_begin_timestamp)); - - if (!begin_ts.isStdOut()) begin_ts << std::dec << app_begin_timestamp.value << std::endl; - - if (begin_ts.fail()) { - rocprofiler::warning("Cannot write to '%s'", begin_ts.name().c_str()); - return; - } + file_plugin_t(void* data) { + if (data) counter_names_ = GetCounterNames(); valid_ = true; } + void WriteHeader(output_type_t type, rocprofiler_tracer_activity_domain_t domain) { + output_file_t* output_file; + switch (domain) { + case ACTIVITY_DOMAIN_HSA_API: { + if (hsa_api_header_written_.load(std::memory_order_relaxed)) return; + output_file = get_output_file(output_type_t::TRACER, ACTIVITY_DOMAIN_HSA_API); + *output_file << "Record_ID,Domain,Function,Start_Timestamp,End_Timestamp,Correlation_ID" + << std::endl; + *output_file << std::endl; + hsa_api_header_written_.exchange(true, std::memory_order_release); + return; + } + case ACTIVITY_DOMAIN_HIP_API: { + if (hip_api_header_written_.load(std::memory_order_relaxed)) return; + output_file = get_output_file(output_type_t::TRACER, ACTIVITY_DOMAIN_HIP_API); + *output_file + << "Record_ID,Domain,Function,Kernel_Name,Start_Timestamp,End_Timestamp,Correlation_ID" + << std::endl; + *output_file << std::endl; + hip_api_header_written_.exchange(true, std::memory_order_release); + return; + } + case ACTIVITY_DOMAIN_ROCTX: { + if (roctx_header_written_.load(std::memory_order_relaxed)) return; + output_file = get_output_file(output_type_t::TRACER, ACTIVITY_DOMAIN_ROCTX); + *output_file << "Record_ID,Domain,ROCTX_ID,Message,Timestamp" << std::endl; + *output_file << std::endl; + roctx_header_written_.exchange(true, std::memory_order_release); + return; + } + case ACTIVITY_DOMAIN_HSA_OPS: { + if (hsa_async_copy_header_written_.load(std::memory_order_relaxed)) return; + output_file = get_output_file(output_type_t::TRACER, ACTIVITY_DOMAIN_HSA_OPS); + *output_file << "Record_ID,Domain,Operation,Start_Timestamp,Stop_Timestamp,Correlation_ID" + << std::endl; + *output_file << std::endl; + hsa_async_copy_header_written_.exchange(true, std::memory_order_release); + return; + } + case ACTIVITY_DOMAIN_HIP_OPS: { + if (hip_activity_header_written_.load(std::memory_order_relaxed)) return; + output_file = get_output_file(output_type_t::TRACER, ACTIVITY_DOMAIN_HIP_OPS); + *output_file << "Record_ID,Domain,Operation,Kernel_Name,Start_Timestamp,Stop_Timestamp," + "Correlation_ID" + << std::endl; + *output_file << std::endl; + hip_activity_header_written_.exchange(true, std::memory_order_release); + return; + } + default: { + if (type == output_type_t::COUNTER) { + if (kernel_dispatches_header_written_.load(std::memory_order_relaxed)) return; + output_file = get_output_file(output_type_t::COUNTER); + + *output_file + << "Dispatch_ID,GPU_ID,Queue_ID,Queue_Index,PID,TID,GRD,WGR,LDS,SCR,Arch_VGPR," + "ACCUM_VGPR,SGPR,Wave_Size,SIG,OBJ,Kernel_Name,Start_Timestamp,End_Timestamp"; + if (counter_names_.size() > 0) { + for (uint32_t i = 0; i < counter_names_.size(); i++) + *output_file << "," << counter_names_[i]; + } + *output_file << std::endl; + *output_file << std::endl; + kernel_dispatches_header_written_.exchange(true, std::memory_order_release); + return; + } else if (type == output_type_t::PC_SAMPLING) { + if (pc_sample_header_written_.load(std::memory_order_relaxed)) return; + output_file = get_output_file(output_type_t::PC_SAMPLING); + *output_file << "Dispatch_ID,Timestamp,GPU_ID,PC_Sample,Shader_Engines" << std::endl; + *output_file << std::endl; + pc_sample_header_written_.exchange(true, std::memory_order_release); + return; + } + return; + } + } + } + std::mutex writing_lock; const char* GetDomainName(rocprofiler_tracer_activity_domain_t domain) { @@ -235,13 +313,21 @@ class file_plugin_t { std::lock_guard lock(writing_lock); if (tracer_record.timestamps.end.value <= 0 && tracer_record.domain != ACTIVITY_DOMAIN_ROCTX) return; + WriteHeader(output_type_t::TRACER, tracer_record.domain); + std::string function_name; std::string kernel_name; std::string roctx_message; uint64_t roctx_id; - if (tracer_record.name) { - if (tracer_record.domain == ACTIVITY_DOMAIN_HIP_API) + if ((tracer_record.operation_id.id == 0 && tracer_record.domain == ACTIVITY_DOMAIN_HIP_OPS)) { + if (tracer_record.name) { kernel_name = rocprofiler::cxx_demangle(tracer_record.name); - if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX) roctx_message = tracer_record.name; + std::string key = "\""; + std::size_t found = kernel_name.rfind(key); + while (found != std::string::npos) { + kernel_name.replace(found, key.length(), "'"); + found = kernel_name.rfind(key, found - 1); + } + } } size_t function_name_size = 0; char* function_name_c = nullptr; @@ -253,6 +339,7 @@ class file_plugin_t { CHECK_ROCPROFILER(rocprofiler_query_hsa_tracer_api_data_info( rocprofiler_session_id_t{0}, ROCPROFILER_HSA_FUNCTION_NAME, tracer_record.api_data_handle, tracer_record.operation_id, &function_name_c)); + if (function_name_c) function_name = std::string(function_name_c); } } if (tracer_record.domain == ACTIVITY_DOMAIN_HIP_API) { @@ -261,51 +348,87 @@ class file_plugin_t { tracer_record.operation_id, &function_name_size)); if (function_name_size > 1) { CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info( - rocprofiler_session_id_t{0}, ROCPROFILER_HIP_FUNCTION_NAME, - tracer_record.api_data_handle, tracer_record.operation_id, &function_name_c)); + session_id, ROCPROFILER_HIP_FUNCTION_NAME, tracer_record.api_data_handle, + tracer_record.operation_id, &function_name_c)); + if (function_name_c) function_name = std::string(function_name_c); } - } - output_file_t* output_file = get_output_file(output_type_t::TRACER, tracer_record.domain); - *output_file << "Record(" << tracer_record.header.id.handle << "), Domain(" - << GetDomainName(tracer_record.domain) << "),"; - if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX && roctx_id >= 0) - *output_file << " ROCTX_ID(" << tracer_record.operation_id.id << "),"; - if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX && tracer_record.name) - *output_file << " ROCTX_Message(" << reinterpret_cast(tracer_record.name) - << "),"; - if (function_name_c) *output_file << " Function(" << function_name_c << "),"; - if (kernel_name.size() > 1) *output_file << " Kernel_Name(" << kernel_name.c_str() << "),"; - if (tracer_record.domain == ACTIVITY_DOMAIN_HSA_OPS || - tracer_record.domain == ACTIVITY_DOMAIN_HIP_OPS) { - switch (tracer_record.operation_id.id) { - case 0: - *output_file << " Operation(DISPATCH_OP),"; - break; - case 1: - *output_file << " Operation(COPY_OP),"; - break; - case 2: - *output_file << " Operation(BARRIER_OP),"; - break; - default: - break; + if (tracer_record.name) { + kernel_name = rocprofiler::cxx_demangle(std::string(tracer_record.name)); + std::string key = "\""; + std::size_t found = kernel_name.rfind(key); + while (found != std::string::npos) { + kernel_name.replace(found, key.length(), "'"); + found = kernel_name.rfind(key, found - 1); + } + // TODO: Change how this API returns a string. } } if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX) { - *output_file << " timestamp(" << tracer_record.timestamps.begin.value << ")"; - } else if (tracer_record.phase == ROCPROFILER_PHASE_EXIT || - tracer_record.phase == ROCPROFILER_PHASE_NONE) { - *output_file << " Begin(" << tracer_record.timestamps.begin.value << "), End(" - << tracer_record.timestamps.end.value << ")"; + if (tracer_record.name) roctx_message = rocprofiler::cxx_demangle(tracer_record.name); + roctx_id = tracer_record.operation_id.id; } - if (tracer_record.domain != ACTIVITY_DOMAIN_ROCTX) - *output_file << ", Correlation_ID(" << tracer_record.correlation_id.value << ")"; - *output_file << '\n'; + char* activity_name = nullptr; + if (tracer_record.domain == ACTIVITY_DOMAIN_HIP_OPS) { + if (tracer_record.api_data_handle.handle) { + kernel_name = rocprofiler::cxx_demangle( + const_cast(reinterpret_cast(tracer_record.api_data_handle.handle))); + } + size_t activity_name_size = 0; + CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info_size( + session_id, ROCPROFILER_HIP_ACTIVITY_NAME, tracer_record.api_data_handle, + tracer_record.operation_id, &activity_name_size)); + if (activity_name_size > 1) { + activity_name = nullptr; + CHECK_ROCPROFILER(rocprofiler_query_hip_tracer_api_data_info( + session_id, ROCPROFILER_HIP_ACTIVITY_NAME, tracer_record.api_data_handle, + tracer_record.operation_id, &activity_name)); + } + } + if (tracer_record.domain == ACTIVITY_DOMAIN_HSA_OPS) { + size_t activity_name_size = 0; + CHECK_ROCPROFILER(rocprofiler_query_hsa_tracer_api_data_info_size( + session_id, ROCPROFILER_HSA_ACTIVITY_NAME, tracer_record.api_data_handle, + tracer_record.operation_id, &activity_name_size)); + if (activity_name_size > 1) { + activity_name = nullptr; + CHECK_ROCPROFILER(rocprofiler_query_hsa_tracer_api_data_info( + session_id, ROCPROFILER_HSA_ACTIVITY_NAME, tracer_record.api_data_handle, + tracer_record.operation_id, &activity_name)); + } + } + // return; + output_file_t* output_file = get_output_file(output_type_t::TRACER, tracer_record.domain); + *output_file << "" << tracer_record.header.id.handle << "," + << GetDomainName(tracer_record.domain); + if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX && roctx_id >= 0) + *output_file << "," << roctx_id; + if (tracer_record.domain == ACTIVITY_DOMAIN_ROCTX) { + if (roctx_message.size() > 1) + *output_file << ",\"" << roctx_message << "\""; + else + *output_file << ","; + } + if (function_name.size() > 1) *output_file << ",\"" << function_name << "\""; + if (activity_name) *output_file << ",\"" << activity_name << "\""; + if (kernel_name.size() > 1) + *output_file << ",\"" << kernel_name.c_str() << "\""; + else if (tracer_record.domain == ACTIVITY_DOMAIN_HIP_API || + tracer_record.domain == ACTIVITY_DOMAIN_HIP_OPS) + *output_file << ","; + if (tracer_record.domain != ACTIVITY_DOMAIN_ROCTX) { + *output_file << "," << tracer_record.timestamps.begin.value << "," + << tracer_record.timestamps.end.value; + *output_file << "," << tracer_record.correlation_id.value; + } else { + *output_file << "," << tracer_record.timestamps.begin.value; + } + *output_file << std::endl; } void FlushProfilerRecord(const rocprofiler_record_profiler_t* profiler_record, rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) { std::lock_guard lock(writing_lock); + WriteHeader(output_type_t::COUNTER, ACTIVITY_DOMAIN_NUMBER); size_t name_length = 0; output_file_t* output_file{nullptr}; output_file = get_output_file(output_type_t::COUNTER); @@ -318,62 +441,42 @@ class file_plugin_t { CHECK_ROCPROFILER(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME, profiler_record->kernel_id, &kernel_name_c)); } - *output_file << std::string("dispatch[") << std::to_string(profiler_record->header.id.handle) - << "], " << std::string("gpu_id(") - << std::to_string(profiler_record->gpu_id.handle) << "), " - << std::string("queue_id(") << std::to_string(profiler_record->queue_id.handle) - << "), " << std::string("queue_index(") - << std::to_string(profiler_record->queue_idx.value) << "), " << std::string("pid(") - << std::to_string(GetPid()) << "), " << std::string("tid(") - << std::to_string(profiler_record->thread_id.value) << ")"; - *output_file << ", " << std::string("grd(") - << std::to_string(profiler_record->kernel_properties.grid_size) << "), " - << std::string("wgr(") - << std::to_string(profiler_record->kernel_properties.workgroup_size) << "), " - << std::string("lds(") + *output_file << std::to_string(profiler_record->header.id.handle) << "," + << std::to_string(profiler_record->gpu_id.handle) << "," + << std::to_string(profiler_record->queue_id.handle) << "," + << std::to_string(profiler_record->queue_idx.value) << "," + << std::to_string(GetPid()) << "," + << std::to_string(profiler_record->thread_id.value); + *output_file << "," << std::to_string(profiler_record->kernel_properties.grid_size) << "," + << std::to_string(profiler_record->kernel_properties.workgroup_size) << "," << std::to_string( ((profiler_record->kernel_properties.lds_size + (lds_block_size - 1)) & ~(lds_block_size - 1))) - << "), " << std::string("scr(") - << std::to_string(profiler_record->kernel_properties.scratch_size) << "), " - << std::string("arch_vgpr(") - << std::to_string(profiler_record->kernel_properties.arch_vgpr_count) << "), " - << std::string("accum_vgpr(") - << std::to_string(profiler_record->kernel_properties.accum_vgpr_count) << "), " - << std::string("sgpr(") - << std::to_string(profiler_record->kernel_properties.sgpr_count) << "), " - << std::string("wave_size(") - << std::to_string(profiler_record->kernel_properties.wave_size) << "), " - << std::string("sig(") + << "," << std::to_string(profiler_record->kernel_properties.scratch_size) << "," + << std::to_string(profiler_record->kernel_properties.arch_vgpr_count) << "," + << std::to_string(profiler_record->kernel_properties.accum_vgpr_count) << "," + << std::to_string(profiler_record->kernel_properties.sgpr_count) << "," + << std::to_string(profiler_record->kernel_properties.wave_size) << "," << std::to_string(profiler_record->kernel_properties.signal_handle); std::string kernel_name = ""; if (name_length > 1) { - kernel_name = rocprofiler::truncate_name(rocprofiler::cxx_demangle(kernel_name_c)); + kernel_name = rocprofiler::cxx_demangle(kernel_name_c); + std::string key = "\""; + std::size_t found = kernel_name.rfind(key); + while (found != std::string::npos) { + kernel_name.replace(found, key.length(), "'"); + found = kernel_name.rfind(key, found - 1); + } } - *output_file << "), " << std::string("obj(") - << std::to_string(profiler_record->kernel_id.handle) << "), " - << std::string("kernel-name(\"") << kernel_name << "\")" - << std::string(", start_time(") - << std::to_string(profiler_record->timestamps.begin.value) << ")" - << std::string(", end_time(") - << std::to_string(profiler_record->timestamps.end.value) << ")"; + *output_file << "," << std::to_string(profiler_record->kernel_id.handle) << ",\"" << kernel_name + << "\"," << std::to_string(profiler_record->timestamps.begin.value) << "," + << std::to_string(profiler_record->timestamps.end.value); // For Counters if (profiler_record->counters) { for (uint64_t i = 0; i < profiler_record->counters_count.value; i++) { if (profiler_record->counters[i].counter_handler.handle > 0) { - size_t counter_name_length = 0; - CHECK_ROCPROFILER(rocprofiler_query_counter_info_size( - session_id, ROCPROFILER_COUNTER_NAME, profiler_record->counters[i].counter_handler, - &counter_name_length)); - if (counter_name_length > 1) { - const char* name_c = nullptr; - CHECK_ROCPROFILER(rocprofiler_query_counter_info( - session_id, ROCPROFILER_COUNTER_NAME, profiler_record->counters[i].counter_handler, - &name_c)); - *output_file << ", " << name_c << " (" - << std::to_string(profiler_record->counters[i].value.value) << ')'; - } + *output_file << "," << std::to_string(profiler_record->counters[i].value.value); } } } @@ -384,14 +487,13 @@ class file_plugin_t { } void FlushPCSamplingRecord(const rocprofiler_record_pc_sample_t* pc_sampling_record) { + WriteHeader(output_type_t::PC_SAMPLING, ACTIVITY_DOMAIN_NUMBER); output_file_t* output_file{nullptr}; output_file = get_output_file(output_type_t::PC_SAMPLING); const auto& sample = pc_sampling_record->pc_sample; - *output_file << "dispatch[" << sample.dispatch_id.value << "], " - << "timestamp(" << sample.timestamp.value << "), " - << "gpu_id(" << sample.gpu_id.handle << "), " - << "pc-sample(" << std::hex << std::showbase << sample.pc << "), " - << "se(" << sample.se << ')' << std::endl; + *output_file << sample.dispatch_id.value << "," << sample.timestamp.value << "," + << sample.gpu_id.handle << "," << std::hex << std::showbase << sample.pc << "," + << sample.se << std::endl; } int WriteBufferRecords(const rocprofiler_record_header_t* begin, const rocprofiler_record_header_t* end, @@ -432,11 +534,17 @@ class file_plugin_t { private: bool valid_{false}; + std::vector counter_names_; - output_file_t roctx_file_{"roctx_trace.txt"}, hsa_api_file_{"hsa_api_trace.txt"}, - hip_api_file_{"hip_api_trace.txt"}, hip_activity_file_{"hcc_ops_trace.txt"}, - hsa_async_copy_file_{"async_copy_trace.txt"}, pc_sample_file_{"pcs_trace.txt"}, - output_file_{"results.txt"}; + std::atomic roctx_header_written_{false}, hsa_api_header_written_{false}, + hip_api_header_written_{false}, hip_activity_header_written_{false}, + hsa_async_copy_header_written_{false}, pc_sample_header_written_{false}, + kernel_dispatches_header_written_{false}; + + output_file_t roctx_file_{"roctx_trace"}, hsa_api_file_{"hsa_api_trace"}, + hip_api_file_{"hip_api_trace"}, hip_activity_file_{"hcc_ops_trace"}, + hsa_async_copy_file_{"async_copy_trace"}, pc_sample_file_{"pcs_trace"}, + output_file_{"results"}; }; file_plugin_t* file_plugin = nullptr; @@ -444,14 +552,15 @@ file_plugin_t* file_plugin = nullptr; } // namespace ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_version, - uint32_t rocprofiler_minor_version) { + uint32_t rocprofiler_minor_version, + void* data) { if (rocprofiler_major_version != ROCPROFILER_VERSION_MAJOR || rocprofiler_minor_version < ROCPROFILER_VERSION_MINOR) return -1; if (file_plugin != nullptr) return -1; - file_plugin = new file_plugin_t(); + file_plugin = new file_plugin_t(data); if (file_plugin->is_valid()) return 0; // The plugin failed to initialized, destroy it and return an error. diff --git a/plugin/perfetto/perfetto.cpp b/plugin/perfetto/perfetto.cpp index 5b36b6cdaa..3d6d2b82ac 100644 --- a/plugin/perfetto/perfetto.cpp +++ b/plugin/perfetto/perfetto.cpp @@ -768,7 +768,7 @@ perfetto_plugin_t* perfetto_plugin = nullptr; } // namespace int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_version, - uint32_t rocprofiler_minor_version) { + uint32_t rocprofiler_minor_version, void* data) { if (rocprofiler_major_version != ROCPROFILER_VERSION_MAJOR || rocprofiler_minor_version > ROCPROFILER_VERSION_MINOR) return -1; diff --git a/src/api/rocprofilerv2.cpp b/src/api/rocprofilerv2.cpp index 59389004fe..9360e22884 100644 --- a/src/api/rocprofilerv2.cpp +++ b/src/api/rocprofilerv2.cpp @@ -357,12 +357,13 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_query_hsa_tracer_api_data_info( // throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_NOT_FOUND); // } // } - if (!(*data = (kind == ROCPROFILER_HSA_FUNCTION_NAME) + *data = (kind == ROCPROFILER_HSA_FUNCTION_NAME) ? rocprofiler::tracer::GetApiCallFunctionName(ACTIVITY_DOMAIN_HSA_API, operation_id) : rocprofiler::GetROCProfilerSingleton() ->GetSession(session_id) ->GetTracer() - ->GetHSAApiDataInfo(kind, api_data_id, operation_id))) + ->GetHSAApiDataInfo(kind, api_data_id, operation_id); + if (*data == nullptr) throw rocprofiler::Exception(ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_INFORMATION_MISSING); API_METHOD_SUFFIX } diff --git a/src/tools/tool.cpp b/src/tools/tool.cpp index e04c70f987..fe26abf034 100644 --- a/src/tools/tool.cpp +++ b/src/tools/tool.cpp @@ -84,6 +84,7 @@ std::atomic session_created{false}; [[maybe_unused]] static rocprofiler_session_id_t session_id; static std::vector filter_ids; static std::vector buffer_ids; +static std::vector counter_names; void warning(const std::string& msg) { std::cerr << msg << std::endl; } @@ -109,7 +110,7 @@ class rocprofiler_plugin_t { if (auto* initialize = reinterpret_cast( dlsym(plugin_handle_, "rocprofiler_plugin_initialize")); initialize != nullptr) - valid_ = initialize(ROCPROFILER_VERSION_MAJOR, ROCPROFILER_VERSION_MINOR) == 0; + valid_ = initialize(ROCPROFILER_VERSION_MAJOR, ROCPROFILER_VERSION_MINOR, &counter_names) == 0; } ~rocprofiler_plugin_t() { @@ -386,7 +387,10 @@ void plugins_load() { if (Dl_info dl_info; dladdr((void*)plugins_load, &dl_info) != 0) { const char* plugin_name = getenv("ROCPROFILER_PLUGIN_LIB"); if (plugin_name == nullptr) { - plugin_name = "libfile_plugin.so"; + if(getenv("OUTPUT_PATH")) + plugin_name = "libfile_plugin.so"; + else + plugin_name = "libcli_plugin.so"; } if (!plugin.emplace(fs::path(dl_info.dli_fname).replace_filename(plugin_name)).is_valid()) { plugin.reset(); @@ -592,6 +596,16 @@ ROCPROFILER_EXPORT bool OnLoad(void* table, uint64_t runtime_version, uint64_t f exit(1); } + std::vector counters = GetCounterNames(); + + if (counters.size() > 0) { + printf("ROCProfilerV2: Collecting the following counters:\n"); + for (size_t i = 0; i < counters.size(); i++) { + counter_names.emplace_back(counters.at(i).c_str()); + printf("- %s\n", counter_names.back()); + } + } + // load the plugins plugins_load(); @@ -606,16 +620,6 @@ ROCPROFILER_EXPORT bool OnLoad(void* table, uint64_t runtime_version, uint64_t f apis_requested.emplace_back(ACTIVITY_DOMAIN_HSA_OPS); if (getenv("ROCPROFILER_ROCTX_TRACE")) apis_requested.emplace_back(ACTIVITY_DOMAIN_ROCTX); - std::vector counters = GetCounterNames(); - std::vector counters_; - - if (counters.size() > 0) { - printf("ROCProfilerV2: Collecting the following counters:\n"); - for (size_t i = 0; i < counters.size(); i++) { - counters_.emplace_back(counters.at(i).c_str()); - printf("- %s\n", counters_.back()); - } - } // ATT Parameters std::vector parameters; std::vector> params; @@ -665,8 +669,8 @@ ROCPROFILER_EXPORT bool OnLoad(void* table, uint64_t runtime_version, uint64_t f rocprofiler_filter_id_t filter_id; [[maybe_unused]] rocprofiler_filter_property_t property = {}; CHECK_ROCPROFILER(rocprofiler_create_filter( - session_id, filter_kind, rocprofiler_filter_data_t{.counters_names = &counters_[0]}, - counters_.size(), &filter_id, property)); + session_id, filter_kind, rocprofiler_filter_data_t{.counters_names = &counter_names[0]}, + counter_names.size(), &filter_id, property)); CHECK_ROCPROFILER(rocprofiler_set_filter_buffer(session_id, filter_id, buffer_id)); filter_ids.emplace_back(filter_id); break; diff --git a/tests-v2/featuretests/profiler/apps/goldentraces/hip_helloworld_golden_traces.txt b/tests-v2/featuretests/profiler/apps/goldentraces/hip_helloworld_golden_traces.txt index 62bd1bb30f..e4a73021db 100755 --- a/tests-v2/featuretests/profiler/apps/goldentraces/hip_helloworld_golden_traces.txt +++ b/tests-v2/featuretests/profiler/apps/goldentraces/hip_helloworld_golden_traces.txt @@ -1,27 +1,4 @@ -0x151a880 agent cpu -0x152b3b0 agent cpu -0x15399c0 agent cpu -0x153b2b0 agent cpu -0x153cb50 agent gpu -0x1541730 agent gpu -0x1593250 agent gpu -0x1597860 agent gpu -0x159be40 agent gpu -0x15a03f0 agent gpu -0x15a49d0 agent gpu -0x15a8f80 agent gpu -5114590922279038 ROCProfilerV2: Collecting the following counters: - GRBM_COUNT Enabling Counter Collection - System minor 0 - System major 9 - agent prop name AMD Instinct MI250X -input string: -GdkknVnqkc - -output string: -HelloWorld -Passed! -dispatch[1], gpu_id(0), queue_id(1), queue_index(1), pid(798749), tid(798749), grd(10), wgr(10), lds(0), scr(0), arch_vgpr(8), accum_vgpr(0), sgpr(16), wave_size(64), sig(139859309564928), obj(1), kernel-name("helloworld"), start_time(5114591230773903), end_time(5114591230779183) -, GRBM_COUNT (11181.000000) \ No newline at end of file +Record_ID(1), GPU_ID(4), Queue_ID(1), Queue_Index(1), Process_ID(1879616), Thread_ID(1879616), Grid_Size(10), Workgroup_Size(10), LDS(0), Scratch_Size(0), Arch_VGPR(8), Accumulative_VGPR(0), SGPR(16), Wave_Size(64), Kernel_Name("helloworld"), Begin_Timestamp(140333224562784), End_Timestamp(140337650854112), GRBM_COUNT(13845.000000) diff --git a/tests-v2/featuretests/profiler/apps/goldentraces/hip_vectoradd_golden_traces.txt b/tests-v2/featuretests/profiler/apps/goldentraces/hip_vectoradd_golden_traces.txt index 44e6faf0a8..8f3907e9b4 100755 --- a/tests-v2/featuretests/profiler/apps/goldentraces/hip_vectoradd_golden_traces.txt +++ b/tests-v2/featuretests/profiler/apps/goldentraces/hip_vectoradd_golden_traces.txt @@ -1,23 +1,4 @@ -0x13615f0 agent cpu -0x1372390 agent cpu -0x13809a0 agent cpu -0x1382290 agent cpu -0x1383b30 agent gpu -0x1388710 agent gpu -0x13da210 agent gpu -0x13de820 agent gpu -0x13e2dc0 agent gpu -0x13e73d0 agent gpu -0x13eb970 agent gpu -0x13eff50 agent gpu -5114627902273635 ROCProfilerV2: Collecting the following counters: - GRBM_COUNT Enabling Counter Collection - System minor 0 - System major 9 - agent prop name AMD Instinct MI250X -hip Device prop succeeded -PASSED! -dispatch[1], gpu_id(0), queue_id(1), queue_index(0), pid(800143), tid(800143), grd(1048576), wgr(256), lds(0), scr(0), arch_vgpr(8), accum_vgpr(0), sgpr(16), wave_size(64), sig(139723878257408), obj(1), kernel-name("vectoradd_float"), start_time(5114628182351915), end_time(5114628182376715) -, GRBM_COUNT (29099.000000) \ No newline at end of file +Record_ID(1), GPU_ID(4), Queue_ID(1), Queue_Index(0), Process_ID(1879580), Thread_ID(1879580), Grid_Size(1048576), Workgroup_Size(256), LDS(0), Scratch_Size(0), Arch_VGPR(8), Accumulative_VGPR(0), SGPR(16), Wave_Size(64), Kernel_Name("vectoradd_float"), Begin_Timestamp(139920974811232), End_Timestamp(139925391173856), GRBM_COUNT(40461.000000) diff --git a/tests-v2/featuretests/profiler/apps/goldentraces/hsa_async_mem_copy_golden_traces.txt b/tests-v2/featuretests/profiler/apps/goldentraces/hsa_async_mem_copy_golden_traces.txt index 623e5261a6..94dd1a87be 100644 --- a/tests-v2/featuretests/profiler/apps/goldentraces/hsa_async_mem_copy_golden_traces.txt +++ b/tests-v2/featuretests/profiler/apps/goldentraces/hsa_async_mem_copy_golden_traces.txt @@ -1,16 +1,12 @@ -0xd1eeb0 agent cpu -0xd4b380 agent gpu -844434431085362 - - -ROCProfiler: Collecting the following counters: +ROCProfilerV2: Collecting the following counters: - GRBM_COUNT - Enabling Counter Collection -Only 1 GPU found with required VRAM. Peer-to-Peer copy will be skipped. -CPU is "AMD Ryzen 9 5950X 16-Core Processor" -GPU1 is "gfx906" +CPU is "AMD Eng Sample: 100-000000248-08_35/21_N" +GPU1 is "gfx90a" +GPU2 is "gfx90a" Copying 4096 bytes from gpu1 memory to system memory... Success! Copying 4096 bytes from system memory to gpu1 memory... Success! +Copying 4096 bytes from gpu1 memory to gpu2 memory... +Success! diff --git a/tests-v2/featuretests/profiler/apps/goldentraces/mpi_vectoradd_golden_traces.txt b/tests-v2/featuretests/profiler/apps/goldentraces/mpi_vectoradd_golden_traces.txt index 1b8bb65acc..bcf8e43d1d 100755 --- a/tests-v2/featuretests/profiler/apps/goldentraces/mpi_vectoradd_golden_traces.txt +++ b/tests-v2/featuretests/profiler/apps/goldentraces/mpi_vectoradd_golden_traces.txt @@ -1,58 +1,17 @@ -0x555942add930 agent cpu -0x555942c51200 agent cpu -0x555942c2aab0 agent cpu -0x555942c2c3a0 agent cpu -0x555942c2dc40 agent gpu -0x555942c327b0 agent gpu -0x555942c864b0 agent gpu -0x555942c8aac0 agent gpu -0x555942c8f0a0 agent gpu -0x555942c93650 agent gpu -0x555942c97c30 agent gpu -0x555942c9c1e0 agent gpu -5114699980260802 ROCProfilerV2: Collecting the following counters: - GRBM_COUNT Enabling Counter Collection -0x22043e0 agent cpu -0x2223990 agent cpu -0x22422f0 agent cpu -0x2243af0 agent cpu -0x2245300 agent gpu -0x2249ee0 agent gpu -0x229ca90 agent gpu -0x22a1070 agent gpu -0x22a5650 agent gpu -0x22a9c00 agent gpu -0x22ae1e0 agent gpu -0x22b2790 agent gpu -5114700383654495 +ROCProfilerV2: Collecting the following counters: +- GRBM_COUNT +Enabling Counter Collection ROCProfilerV2: Collecting the following counters: - GRBM_COUNT Enabling Counter Collection device count and rank is8: 2 Rank Id: 0 | Device Id : 0 | Num Devices: 8 -0x163d3f0 agent cpu -0x165c9c0 agent cpu -0x167b320 agent cpu -0x167cb20 agent cpu -0x167e330 agent gpu -0x1682f10 agent gpu -0x16d5ac0 agent gpu -0x16da0a0 agent gpu -0x16de680 agent gpu -0x16e2c30 agent gpu -0x16e7210 agent gpu -0x16eb7c0 agent gpu -5114700398718649 -ROCProfilerV2: Collecting the following counters: -- GRBM_COUNT -Enabling Counter Collection device count and rank is8: 2 Rank Id: 1 | Device Id : 1 | Num Devices: 8 Max error: 0.000000 Max error: 0.000000 -dispatch[1], gpu_id(0), queue_id(1), queue_index(0), pid(800362), tid(800362), grd(1048576), wgr(256), lds(0), scr(0), arch_vgpr(12), accum_vgpr(4), sgpr(16), wave_size(64), sig(140592265663488), obj(1), kernel-name("add"), start_time(5114700670020244), end_time(5114700670397685) -, GRBM_COUNT (497443.000000) -dispatch[1], gpu_id(1), queue_id(1), queue_index(0), pid(800363), tid(800363), grd(1048576), wgr(256), lds(0), scr(0), arch_vgpr(12), accum_vgpr(4), sgpr(16), wave_size(64), sig(139748397077504), obj(1), kernel-name("add"), start_time(5114700683506777), end_time(5114700683876218) -, GRBM_COUNT (485724.000000) \ No newline at end of file +Record_ID(1), GPU_ID(4), Queue_ID(1), Queue_Index(0), Process_ID(1879535), Thread_ID(1879535), Grid_Size(1048576), Workgroup_Size(256), LDS(0), Scratch_Size(0), Arch_VGPR(12), Accumulative_VGPR(4), SGPR(32), Wave_Size(64), Kernel_Name("add"), Begin_Timestamp(140358658822704), End_Timestamp(140358853089512), GRBM_COUNT(505731.000000) +Record_ID(1), GPU_ID(5), Queue_ID(1), Queue_Index(0), Process_ID(1879536), Thread_ID(1879536), Grid_Size(1048576), Workgroup_Size(256), LDS(0), Scratch_Size(0), Arch_VGPR(12), Accumulative_VGPR(4), SGPR(32), Wave_Size(64), Kernel_Name("add"), Begin_Timestamp(140518780572208), End_Timestamp(140518958147816), GRBM_COUNT(496299.000000) diff --git a/tests-v2/featuretests/profiler/profiler_gtest.cpp b/tests-v2/featuretests/profiler/profiler_gtest.cpp index 5e48082a68..76f2d71af4 100644 --- a/tests-v2/featuretests/profiler/profiler_gtest.cpp +++ b/tests-v2/featuretests/profiler/profiler_gtest.cpp @@ -27,13 +27,13 @@ THE SOFTWARE. #include "rocprofiler.h" #include +#include #include #include #include #include #include "src/utils/helper.h" -#include "utils/test_utils.h" #include "utils/csv_parser.h" #include "src/utils/logger.h" #include "apps/hip_kernels.h" @@ -106,47 +106,18 @@ void ApplicationParser::SetApplicationEnv(const char* app_name) { * and saves them in a vector. */ void ApplicationParser::GetKernelInfoForRunningApplication( - std::vector* kernel_info_output) { - KernelInfo kinfo; + std::vector* kernel_info_output) { + profiler_kernel_info_t kinfo; for (std::string line : output_lines) { - if (std::regex_match(line, std::regex("(dispatch)(.*)"))) { - int spos = line.find("["); - int epos = line.find("]", spos); - std::string sub = line.substr(spos + 1, epos - spos - 1); - kinfo.dispatch_id = sub; - kernel_info_output->push_back(kinfo); - - // Kernel-Name - size_t found = line.find("kernel-name"); - if (found != std::string::npos) { - int spos = found; - int epos = line.find(")", spos); - int length = std::string("kernel-name").length(); - std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); - - kinfo.kernel_name = sub; - kernel_info_output->push_back(kinfo); - } - // Start-Time - found = line.find("start_time"); - if (found != std::string::npos) { - int spos = found; - int epos = line.find(",", spos); - int length = std::string("start_time").length(); - std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); - kinfo.start_time = sub; - kernel_info_output->push_back(kinfo); - } - // End-Time - found = line.find("end_time"); - if (found != std::string::npos) { - int spos = line.find(",", found); - int epos = line.find(")", spos); - std::string sub = line.substr(spos + 1, epos - spos - 1); - kinfo.end_time = sub; - kernel_info_output->push_back(kinfo); - } + // Skip all the lines until "Record_ID" is found + if (line.empty() || line.find("Record_ID") == std::string::npos) { + continue; // Skip to the next line if "Record_ID" is found } + + // Parse individual values and store them in the dispatch struct + tokenize_profiler_output(line, kinfo); + + kernel_info_output->push_back(kinfo); } } @@ -154,8 +125,9 @@ void ApplicationParser::GetKernelInfoForRunningApplication( * Parses kernel-names from a pre-saved golden out files * and saves them in a vector. */ -void ApplicationParser::GetKernelInfoForGoldenOutput(const char* app_name, std::string file_name, - std::vector* kernel_info_output) { +void ApplicationParser::GetKernelInfoForGoldenOutput( + const char* app_name, std::string file_name, + std::vector* kernel_info_output) { std::string entry; std::string path = GetRunningPath(running_path); entry = path.append(golden_trace_path) + file_name; @@ -195,52 +167,21 @@ void ApplicationParser::ProcessApplication(std::stringstream& ss) { * Parses kernel-info for golden output file * and saves them in a vector. */ -void ApplicationParser::ParseKernelInfoFields(const std::string& s, - std::vector* kernel_info_output) { +void ApplicationParser::ParseKernelInfoFields( + const std::string& s, std::vector* kernel_info_output) { std::string line; - KernelInfo kinfo; + profiler_kernel_info_t kinfo; std::ifstream golden_file(s); while (!golden_file.eof()) { getline(golden_file, line); - if (std::regex_match(line, std::regex("(dispatch)(.*)"))) { - int spos = line.find("["); - int epos = line.find("]", spos); - std::string sub = line.substr(spos + 1, epos - spos - 1); - kinfo.dispatch_id = sub; - kernel_info_output->push_back(kinfo); - - // Kernel-Name - size_t found = line.find("kernel-name"); - if (found != std::string::npos) { - int spos = found; - int epos = line.find(")", spos); - int length = std::string("kernel-name").length(); - std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); - - kinfo.kernel_name = sub; - kernel_info_output->push_back(kinfo); - } - // Start-Time - found = line.find("start_time"); - if (found != std::string::npos) { - int spos = found; - int epos = line.find(",", spos); - int length = std::string("start_time").length(); - std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); - kinfo.start_time = sub; - kernel_info_output->push_back(kinfo); - } - // End-Time - found = line.find("end_time"); - if (found != std::string::npos) { - int spos = line.find(",", found); - int epos = line.find(")", spos); - std::string sub = line.substr(spos + 1, epos - spos - 1); - kinfo.end_time = sub; - kernel_info_output->push_back(kinfo); - } + // Skip all the lines until "Record_ID" is found + if (line.empty() || line.find("Record_ID") == std::string::npos) { + continue; // Skip to the next line if "Record_ID" is found } + // Parse individual values and store them in the dispatch struct + tokenize_profiler_output(line, kinfo); + kernel_info_output->push_back(kinfo); } golden_file.close(); } @@ -260,7 +201,7 @@ constexpr auto kGoldenOutputMpi = "mpi_vectoradd_golden_traces.txt"; class HelloWorldTest : public ProfilerTest { protected: - std::vector golden_kernel_info; + std::vector golden_kernel_info; void SetUp() { ProfilerTest::SetUp("hip_helloworld"); GetKernelInfoForGoldenOutput("hip_helloworld", kGoldenOutputHelloworld, &golden_kernel_info); @@ -271,7 +212,7 @@ class HelloWorldTest : public ProfilerTest { // profiler output TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { // kernel info in current profiler run - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); @@ -283,20 +224,19 @@ TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGolde // profiler output TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenKernelNamessMatchWithGoldenOutput) { // kernel info in current profiler run - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); - + ASSERT_TRUE(golden_kernel_info.size()); EXPECT_EQ(golden_kernel_info[0].kernel_name, current_kernel_info[0].kernel_name); - EXPECT_EQ(golden_kernel_info[1].kernel_name, current_kernel_info[1].kernel_name); } // Test:3 Compares order of kernel-names in golden output against current // profiler output TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenKernelDurationShouldBePositive) { // kernel info in current profiler run - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); @@ -308,14 +248,14 @@ TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenKernelDurationShouldBePosit // profiler output TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenEndTimeIsGreaterThenStartTime) { // kernel info in current profiler run - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); for (auto& itr : current_kernel_info) { - if (!(itr.start_time).empty() && !(itr.end_time).empty()) { - EXPECT_GT(itr.end_time, itr.start_time); + if (!(itr.begin_time).empty() && !(itr.end_time).empty()) { + EXPECT_GT(itr.end_time, itr.begin_time); } } } @@ -328,7 +268,7 @@ TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenEndTimeIsGreaterThenStartTi class VectorAddTest : public ProfilerTest { protected: - std::vector golden_kernel_info; + std::vector golden_kernel_info; void SetUp() { ProfilerTest::SetUp("hip_vectoradd"); GetKernelInfoForGoldenOutput("hip_vectoradd", kGoldenOutputVectorAdd, &golden_kernel_info); @@ -338,7 +278,7 @@ class VectorAddTest : public ProfilerTest { // Test:1 Compares total num of kernel-names in golden output against current // profiler output TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); @@ -349,20 +289,20 @@ TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGolden // Test:2 Compares order of kernel-names in golden output against current // profiler output TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenKernelNamessMatchWithGoldenOutput) { - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); + ASSERT_TRUE(golden_kernel_info.size()); EXPECT_EQ(golden_kernel_info[0].kernel_name, current_kernel_info[0].kernel_name); - EXPECT_EQ(golden_kernel_info[1].kernel_name, current_kernel_info[1].kernel_name); } // Test:3 Compares order of kernel-names in golden output against current // profiler output TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenKernelDurationShouldBePositive) { // kernel info in current profiler run - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); @@ -374,14 +314,14 @@ TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenKernelDurationShouldBePositi // profiler output TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenEndTimeIsGreaterThenStartTime) { // kernel info in current profiler run - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); for (auto& itr : current_kernel_info) { - if (!(itr.start_time).empty() && !(itr.end_time).empty()) { - EXPECT_GT(itr.end_time, itr.start_time); + if (!(itr.begin_time).empty() && !(itr.end_time).empty()) { + EXPECT_GT(itr.end_time, itr.begin_time); } } } @@ -394,7 +334,7 @@ TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenEndTimeIsGreaterThenStartTim class HSATest : public ProfilerTest { protected: - std::vector golden_kernel_info; + std::vector golden_kernel_info; void SetUp() { ProfilerTest::SetUp("hsa_async_mem_copy"); GetKernelInfoForGoldenOutput("hsa_async_mem_copy", kGOldenOutputAsyncCopy, &golden_kernel_info); @@ -405,7 +345,7 @@ class HSATest : public ProfilerTest { // we dont collect any counters by default. Expectation is, both vectors are // empty TEST_F(HSATest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); @@ -423,7 +363,7 @@ TEST_F(HSATest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput class OpenMPTest : public ProfilerTest { protected: - std::vector golden_kernel_info; + std::vector golden_kernel_info; void SetUp() { ProfilerTest::SetUp("openmp_helloworld"); GetKernelInfoForGoldenOutput("openmp_helloworld", kGoldenOutputOpenMP, &golden_kernel_info); @@ -433,7 +373,7 @@ class OpenMPTest : public ProfilerTest { // Test:1 Compares total num of kernel-names in golden output against current // profiler output TEST_F(OpenMPTest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); @@ -444,7 +384,7 @@ TEST_F(OpenMPTest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOut // Test:2 Compares order of kernel-names in golden output against current // profiler output TEST_F(OpenMPTest, WhenRunningProfilerWithAppThenKernelNamessMatchWithGoldenOutput) { - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); @@ -457,7 +397,7 @@ TEST_F(OpenMPTest, WhenRunningProfilerWithAppThenKernelNamessMatchWithGoldenOutp // profiler output TEST_F(OpenMPTest, WhenRunningProfilerWithAppThenKernelDurationShouldBePositive) { // kernel info in current profiler run - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); @@ -469,14 +409,14 @@ TEST_F(OpenMPTest, WhenRunningProfilerWithAppThenKernelDurationShouldBePositive) // profiler output TEST_F(OpenMPTest, WhenRunningProfilerWithAppThenEndTimeIsGreaterThenStartTime) { // kernel info in current profiler run - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); for (auto& itr : current_kernel_info) { if (!(itr.end_time).empty()) { - EXPECT_GT(itr.end_time, itr.start_time); + EXPECT_GT(itr.end_time, itr.begin_time); } } } @@ -543,7 +483,7 @@ void MPITest::ExecuteAndParseApplication(std::stringstream& ss) { // Test:1 Compares total num of kernel-names in golden output against current // profiler output TEST_F(MPITest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); @@ -770,13 +710,13 @@ class ProfilerAPITest : public ::testing::Test { CheckApi(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME, profiler_record->kernel_id, &kernel_name_c)); // int gpu_index = profiler_record->gpu_id.handle; - // uint64_t start_time = profiler_record->timestamps.begin.value; + // uint64_t begin_time = profiler_record->timestamps.begin.value; // uint64_t end_time = profiler_record->timestamps.end.value; // printf( // "Kernel Info:\n\tGPU Index: %d\n\tKernel Name: %s\n\tStart " // "Time: " // "%lu\n\tEnd Time: %lu\n", - // gpu_index, kernel_name_c, start_time, end_time); + // gpu_index, kernel_name_c, begin_time, end_time); } CheckApi(rocprofiler_next_record(record, &record, session_id, buffer_id)); } @@ -1173,33 +1113,34 @@ TEST(ProfilerMPTest, WhenRunningMultiProcessTestItPasses) { // bool hasFile() { return hasFileInDir("file_test_name", "."); } // }; -// TEST_F(VectorAddFileOnlyTest, WhenRunningProfilerWithFilePluginTest) { EXPECT_EQ(hasFile(), true); } +// TEST_F(VectorAddFileOnlyTest, WhenRunningProfilerWithFilePluginTest) { EXPECT_EQ(hasFile(), +// true); } // class VectorAddFolderOnlyTest : public FilePluginTest { // protected: // virtual void SetUp() { -// RunApplication("hip_vectoradd", " --hsa-activity --hip-activity -d ./plugin_test_folder_path"); +// RunApplication("hip_vectoradd", " --hsa-activity --hip-activity -d +// ./plugin_test_folder_path"); // } // virtual void TearDown() { -// std::experimental::filesystem::remove_all("./plugin_test_folder_path"); -// } -// bool hasFile() { return hasFileInDir("", "./plugin_test_folder_path"); } +// std::experimental::filesystem::remove_all("./plugin_test_folder_path"); +// } bool hasFile(){ return hasFileInDir("", "./plugin_test_folder_path"); } // }; // TEST_F(VectorAddFolderOnlyTest, WhenRunningProfilerWithFilePluginTest) { // EXPECT_EQ(hasFile(), true); // } + // class VectorAddFileAndFolderTest : public FilePluginTest { // protected: // virtual void SetUp() { -// RunApplication("hip_vectoradd", -// " --hip-activity -d ./plugin_test_folder_path -o file_test_name"); +// RunApplication("hip_vectoradd", " --hip-activity -d ./plugin_test_folder_path -o +// file_test_name"); // } // virtual void TearDown() { -// std::experimental::filesystem::remove_all("./plugin_test_folder_path"); -// } -// bool hasFile() { return hasFileInDir("file_test_name", "./plugin_test_folder_path"); } +// std::experimental::filesystem::remove_all("./plugin_test_folder_path"); +// } bool hasFile(){ return hasFileInDir("file_test_name", "./plugin_test_folder_path"); } // }; // TEST_F(VectorAddFileAndFolderTest, WhenRunningProfilerWithFilePluginTest) { @@ -1210,7 +1151,8 @@ TEST(ProfilerMPTest, WhenRunningMultiProcessTestItPasses) { // protected: // virtual void SetUp() { // setenv("MPI_RANK", "7", true); -// RunApplication("hip_vectoradd", " --hip-activity -d ./plugin_test_folder_path -o test_%rank_"); +// RunApplication("hip_vectoradd", " --hip-activity -d ./plugin_test_folder_path -o +// test_%rank_"); // } // virtual void TearDown() { // std::experimental::filesystem::remove_all("./plugin_test_folder_path"); @@ -1238,14 +1180,14 @@ TEST(ProfilerMPTest, WhenRunningMultiProcessTestItPasses) { // protected: // virtual void SetUp() { // setenv("MPI_RANK", "7", true); -// RunApplication("hip_vectoradd", -// " -d ./plugin_test_folder_path -o test_%rank_ --plugin perfetto"); +// RunApplication("hip_vectoradd", " -d ./plugin_test_folder_path -o test_%rank_ --plugin +// perfetto"); // } // virtual void TearDown() { // std::experimental::filesystem::remove_all("./plugin_test_folder_path"); // unsetenv("MPI_RANK"); // } -// bool hasFile() { return hasFileInDir("test_7_", "./plugin_test_folder_path"); } +// bool hasFile(){ return hasFileInDir("test_7_", "./plugin_test_folder_path"); } // }; // TEST_F(VectorAddPerfettoMPITest, WhenRunningProfilerWithPerfettoTest) { diff --git a/tests-v2/featuretests/profiler/profiler_gtest.h b/tests-v2/featuretests/profiler/profiler_gtest.h index 5d526ed185..7161b5a2e5 100644 --- a/tests-v2/featuretests/profiler/profiler_gtest.h +++ b/tests-v2/featuretests/profiler/profiler_gtest.h @@ -40,6 +40,8 @@ THE SOFTWARE. #include #include +#include "utils/test_utils.h" + /* --------------------------------------------------------------------------*/ /** * @Synopsis Implementation of a Parser class for Profiler output @@ -54,19 +56,6 @@ class ApplicationParser : public ::testing::Test { protected: virtual void SetUp(const char* app_name) { SetApplicationEnv(app_name); } virtual void TearDown() {} - //!< This can be appended for other kernel info fields; eg: Agent-Name etc. - struct KernelInfo { - std::string dispatch_id; - std::string gpu_id; - std::string queue_id; - std::string queue_index; - std::string pid; - std::string tid; - std::string obj; - std::string kernel_name; - std::string start_time; - std::string end_time; - }; //!< saves lines of profiler output std::vector output_lines; @@ -78,11 +67,11 @@ class ApplicationParser : public ::testing::Test { //!< Parses kernel-info from a pre-saved golden out files // and saves them in a vector. void GetKernelInfoForGoldenOutput(const char* app_name, std::string filename, - std::vector* kernel_info_output); + std::vector* kernel_info_output); //!< Parses kernel-info after running profiler against curent application // and saves them in a vector. - void GetKernelInfoForRunningApplication(std::vector* kernel_info_output); + void GetKernelInfoForRunningApplication(std::vector* kernel_info_output); private: //!< Runs a given appllication and saves profiler output. @@ -92,7 +81,7 @@ class ApplicationParser : public ::testing::Test { //!< Parses kernel info fields from given input // i.e: kernel_names, kernel_duration - void ParseKernelInfoFields(const std::string& s, std::vector* kernel_info_output); + void ParseKernelInfoFields(const std::string& s, std::vector* kernel_info_output); }; /* --------------------------------------------------------------------------*/ diff --git a/tests-v2/featuretests/tracer/apps/goldentraces/hip_helloworld_golden_traces.txt b/tests-v2/featuretests/tracer/apps/goldentraces/hip_helloworld_golden_traces.txt index f1056fb069..0e6f5e1679 100755 --- a/tests-v2/featuretests/tracer/apps/goldentraces/hip_helloworld_golden_traces.txt +++ b/tests-v2/featuretests/tracer/apps/goldentraces/hip_helloworld_golden_traces.txt @@ -1,14 +1,11 @@ -0x1fd8190 agent cpu -0x1fd8ee0 agent gpu -1124305956573108 Enabling API Tracing -Record(2), Domain(HIP_API_DOMAIN), Function(hipGetDeviceProperties), Begin(140736506355256), End(1124305958759521), Correlation_ID(1) -Record(4), Domain(HIP_API_DOMAIN), Function(hipMalloc), Begin(140736506355256), End(1124305958873631), Correlation_ID(2) -Record(6), Domain(HIP_API_DOMAIN), Function(hipMalloc), Begin(140736506355256), End(1124305958880271), Correlation_ID(3) -Record(8), Domain(HIP_API_DOMAIN), Function(hipMemcpy), Begin(140736506355256), End(1124306216889435), Correlation_ID(4) -Record(10), Domain(HIP_API_DOMAIN), Function(__hipPushCallConfiguration), Begin(140736506355256), End(1124306216895505), Correlation_ID(5) -Record(12), Domain(HIP_API_DOMAIN), Function(__hipPopCallConfiguration), Begin(140736506355256), End(1124306216897125), Correlation_ID(6) -Record(14), Domain(HIP_API_DOMAIN), Function(hipLaunchKernel), Kernel_Name(helloworld(char*, char*)), Begin(140736506355240), End(1124306217234255), Correlation_ID(7) -Record(16), Domain(HIP_API_DOMAIN), Function(hipMemcpy), Begin(140736506355256), End(1124306217720846), Correlation_ID(8) -Record(18), Domain(HIP_API_DOMAIN), Function(hipFree), Begin(140736506355272), End(1124306217730836), Correlation_ID(9) -Record(20), Domain(HIP_API_DOMAIN), Function(hipFree), Begin(140736506355272), End(1124306217734156), Correlation_ID(10) \ No newline at end of file +Record_ID(2), Domain(HIP_API_DOMAIN), Function(hipGetDeviceProperties), Start_Timestamp(611510516147389), End_Timestamp(611510516154633), Correlation_ID(1) +Record_ID(4), Domain(HIP_API_DOMAIN), Function(hipMalloc), Start_Timestamp(611510516160324), End_Timestamp(611510516183168), Correlation_ID(2) +Record_ID(6), Domain(HIP_API_DOMAIN), Function(hipMalloc), Start_Timestamp(611510516183599), End_Timestamp(611510516189250), Correlation_ID(3) +Record_ID(8), Domain(HIP_API_DOMAIN), Function(hipMemcpy), Start_Timestamp(611510516192225), End_Timestamp(611510784335292), Correlation_ID(4) +Record_ID(10), Domain(HIP_API_DOMAIN), Function(__hipPushCallConfiguration), Start_Timestamp(611510784338909), End_Timestamp(611510784341915), Correlation_ID(5) +Record_ID(12), Domain(HIP_API_DOMAIN), Function(__hipPopCallConfiguration), Start_Timestamp(611510784343228), End_Timestamp(611510784343538), Correlation_ID(6) +Record_ID(14), Domain(HIP_API_DOMAIN), Function(hipLaunchKernel), Kernel_Name(helloworld(char*, char*)), Start_Timestamp(611510784351053), End_Timestamp(611510784872311), Correlation_ID(7) +Record_ID(16), Domain(HIP_API_DOMAIN), Function(hipMemcpy), Start_Timestamp(611510784873914), End_Timestamp(611510784891167), Correlation_ID(8) +Record_ID(18), Domain(HIP_API_DOMAIN), Function(hipFree), Start_Timestamp(611510784893482), End_Timestamp(611510784902690), Correlation_ID(9) +Record_ID(20), Domain(HIP_API_DOMAIN), Function(hipFree), Start_Timestamp(611510784904944), End_Timestamp(611510784907128), Correlation_ID(10) diff --git a/tests-v2/featuretests/tracer/apps/goldentraces/hsa_api_async_trace.txt b/tests-v2/featuretests/tracer/apps/goldentraces/hsa_api_async_trace.txt index ac3c49a08b..cef7c22cd7 100644 --- a/tests-v2/featuretests/tracer/apps/goldentraces/hsa_api_async_trace.txt +++ b/tests-v2/featuretests/tracer/apps/goldentraces/hsa_api_async_trace.txt @@ -1,55 +1,80 @@ -0x22c6900 agent cpu -0x22c8bf0 agent gpu -1125004784725841 Enabling API Tracing -Only 1 GPU found with required VRAM. Peer-to-Peer copy will be skipped. -CPU is "AMD Ryzen 9 3900X 12-Core Processor" -GPU1 is "gfx900" +CPU is "AMD Eng Sample: 100-000000248-08_35/21_N" +GPU1 is "gfx90a" +GPU2 is "gfx90a" Copying 4096 bytes from gpu1 memory to system memory... Success! Copying 4096 bytes from system memory to gpu1 memory... Success! -Record(3), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Begin(140734995053648), End(1125004785031455), Correlation_ID(2) -Record(6), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Begin(140734995053280), End(1125004785035995), Correlation_ID(4) -Record(8), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Begin(140734995053280), End(1125004785036555), Correlation_ID(5) -Record(10), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Begin(140734995053280), End(1125004785037085), Correlation_ID(6) -Record(11), Domain(HSA_API_DOMAIN), Function(hsa_amd_agent_iterate_memory_pools), Begin(140734995053648), End(1125004785037355), Correlation_ID(3) -Record(12), Domain(HSA_API_DOMAIN), Function(hsa_iterate_agents), Begin(140734995054064), End(1125004785037625), Correlation_ID(1) -Record(15), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Begin(140734995053632), End(1125004785038395), Correlation_ID(8) -Record(17), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Begin(140734995053632), End(1125004785038995), Correlation_ID(9) -Record(20), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Begin(140734995053248), End(1125004785043495), Correlation_ID(11) -Record(22), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Begin(140734995053248), End(1125004785043975), Correlation_ID(12) -Record(24), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Begin(140734995053248), End(1125004785044445), Correlation_ID(13) -Record(25), Domain(HSA_API_DOMAIN), Function(hsa_amd_agent_iterate_memory_pools), Begin(140734995053632), End(1125004785044685), Correlation_ID(10) -Record(26), Domain(HSA_API_DOMAIN), Function(hsa_iterate_agents), Begin(140734995054064), End(1125004785044935), Correlation_ID(7) -Record(28), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_allocate), Begin(140734995054048), End(1125004785087865), Correlation_ID(14) -Record(30), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_allocate), Begin(140734995054048), End(1125004785211957), Correlation_ID(15) -Record(32), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Begin(140734995054048), End(1125004785213057), Correlation_ID(16) -Record(34), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Begin(140734995054048), End(1125004785219267), Correlation_ID(17) -Record(36), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_fill), Begin(140734995053904), End(1125004785227097), Correlation_ID(18) -Record(38), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_fill), Begin(140734995053904), End(1125004798095315), Correlation_ID(19) -Record(40), Domain(HSA_API_DOMAIN), Function(hsa_amd_agents_allow_access), Begin(140734995053904), End(1125004798155806), Correlation_ID(20) -Record(42), Domain(HSA_API_DOMAIN), Function(hsa_signal_create), Begin(140734995053904), End(1125004798159096), Correlation_ID(21) -Record(44), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_async_copy), Begin(140734995053872), End(1125004798912335), Correlation_ID(22) -Record(47), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_copy_engine_status), Begin(140734995053904), End(1125004799939028), Correlation_ID(23) -Record(49), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_async_copy_on_engine), Begin(140734995053856), End(1125004799950448), Correlation_ID(24) -Record(51), Domain(HSA_API_DOMAIN), Function(hsa_signal_wait_relaxed), Begin(140734995053888), End(1125004799952668), Correlation_ID(25) -Record(54), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_fill), Begin(140734995053904), End(1125004799996599), Correlation_ID(26) -Record(56), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_fill), Begin(140734995053904), End(1125004800133540), Correlation_ID(27) -Record(58), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_fill), Begin(140734995053904), End(1125004800146430), Correlation_ID(28) -Record(60), Domain(HSA_API_DOMAIN), Function(hsa_amd_agents_allow_access), Begin(140734995053904), End(1125004800150920), Correlation_ID(29) -Record(62), Domain(HSA_API_DOMAIN), Function(hsa_signal_create), Begin(140734995053904), End(1125004800152290), Correlation_ID(30) -Record(64), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_async_copy), Begin(140734995053872), End(1125004800157800), Correlation_ID(31) -Record(66), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_copy_engine_status), Begin(140734995053904), End(1125004800158741), Correlation_ID(32) -Record(68), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_async_copy_on_engine), Begin(140734995053856), End(1125004800162091), Correlation_ID(33) -Record(71), Domain(HSA_API_DOMAIN), Function(hsa_signal_wait_relaxed), Begin(140734995053888), End(1125004800172941), Correlation_ID(34) -Record(73), Domain(HSA_API_DOMAIN), Function(hsa_signal_store_screlease), Begin(140734995053920), End(1125004800174791), Correlation_ID(35) -Record(75), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_async_copy), Begin(140734995053872), End(1125004800181531), Correlation_ID(36) -Record(78), Domain(HSA_API_DOMAIN), Function(hsa_signal_wait_relaxed), Begin(140734995053888), End(1125004800184931), Correlation_ID(37) -Record(81), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_free), Begin(140734995054064), End(1125004800235281), Correlation_ID(38) -Record(83), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_free), Begin(140734995054064), End(1125004800238011), Correlation_ID(39) -Record(46), Domain(HSA_OPS_DOMAIN), Operation(COPY_OP), Begin(1125004798940727), End(1125004798945616), Correlation_ID(22) -Record(53), Domain(HSA_OPS_DOMAIN), Operation(COPY_OP), Begin(1125004799956748), End(1125004799961636), Correlation_ID(24) -Record(70), Domain(HSA_OPS_DOMAIN), Operation(COPY_OP), Begin(1125004800167536), End(1125004800171091), Correlation_ID(31) -Record(77), Domain(HSA_OPS_DOMAIN), Operation(COPY_OP), Begin(1125004800175831), End(1125004800179534), Correlation_ID(33) -Record(80), Domain(HSA_OPS_DOMAIN), Operation(COPY_OP), Begin(1125004800192274), End(1125004800195829), Correlation_ID(36) +Copying 4096 bytes from gpu1 memory to gpu2 memory... +Success! +Record_ID(68), Domain(HSA_OPS_DOMAIN), Operation_Name(COPY), Start_Timestamp(611575185146090), End_Timestamp(611575185153450), Correlation_ID(33) +Record_ID(91), Domain(HSA_OPS_DOMAIN), Operation_Name(COPY), Start_Timestamp(611575186828236), End_Timestamp(611575186834316), Correlation_ID(42) +Record_ID(98), Domain(HSA_OPS_DOMAIN), Operation_Name(COPY), Start_Timestamp(611575186839756), End_Timestamp(611575186842796), Correlation_ID(44) +Record_ID(101), Domain(HSA_OPS_DOMAIN), Operation_Name(COPY), Start_Timestamp(611575186845596), End_Timestamp(611575186848796), Correlation_ID(47) +Record_ID(118), Domain(HSA_OPS_DOMAIN), Operation_Name(COPY), Start_Timestamp(611575191866832), End_Timestamp(611575192127791), Correlation_ID(54) +Record_ID(127), Domain(HSA_OPS_DOMAIN), Operation_Name(COPY), Start_Timestamp(611575193184345), End_Timestamp(611575193190105), Correlation_ID(60) +Record_ID(3), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Start_Timestamp(611575178942576), End_Timestamp(611575178943147), Correlation_ID(2) +Record_ID(6), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Start_Timestamp(611575178951443), End_Timestamp(611575178951814), Correlation_ID(4) +Record_ID(8), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Start_Timestamp(611575178952044), End_Timestamp(611575178952345), Correlation_ID(5) +Record_ID(10), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Start_Timestamp(611575178952605), End_Timestamp(611575178952776), Correlation_ID(6) +Record_ID(11), Domain(HSA_API_DOMAIN), Function(hsa_amd_agent_iterate_memory_pools), Start_Timestamp(611575178949880), End_Timestamp(611575178952996), Correlation_ID(3) +Record_ID(12), Domain(HSA_API_DOMAIN), Function(hsa_iterate_agents), Start_Timestamp(611575178937616), End_Timestamp(611575178953207), Correlation_ID(1) +Record_ID(15), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Start_Timestamp(611575178953637), End_Timestamp(611575178953808), Correlation_ID(8) +Record_ID(17), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Start_Timestamp(611575178954088), End_Timestamp(611575178954239), Correlation_ID(9) +Record_ID(19), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Start_Timestamp(611575178954419), End_Timestamp(611575178956403), Correlation_ID(10) +Record_ID(21), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Start_Timestamp(611575178956593), End_Timestamp(611575178956733), Correlation_ID(11) +Record_ID(23), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Start_Timestamp(611575178956964), End_Timestamp(611575178957204), Correlation_ID(12) +Record_ID(26), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Start_Timestamp(611575178957695), End_Timestamp(611575178957876), Correlation_ID(14) +Record_ID(28), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Start_Timestamp(611575178958066), End_Timestamp(611575178958216), Correlation_ID(15) +Record_ID(30), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Start_Timestamp(611575178958427), End_Timestamp(611575178958557), Correlation_ID(16) +Record_ID(31), Domain(HSA_API_DOMAIN), Function(hsa_amd_agent_iterate_memory_pools), Start_Timestamp(611575178957425), End_Timestamp(611575178958727), Correlation_ID(13) +Record_ID(33), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Start_Timestamp(611575178958918), End_Timestamp(611575178959058), Correlation_ID(17) +Record_ID(36), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Start_Timestamp(611575178959469), End_Timestamp(611575178959609), Correlation_ID(19) +Record_ID(38), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Start_Timestamp(611575178960731), End_Timestamp(611575178960871), Correlation_ID(20) +Record_ID(40), Domain(HSA_API_DOMAIN), Function(hsa_amd_agent_memory_pool_get_info), Start_Timestamp(611575178961893), End_Timestamp(611575178962164), Correlation_ID(21) +Record_ID(42), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_get_info), Start_Timestamp(611575178962354), End_Timestamp(611575178962515), Correlation_ID(22) +Record_ID(43), Domain(HSA_API_DOMAIN), Function(hsa_amd_agent_iterate_memory_pools), Start_Timestamp(611575178959268), End_Timestamp(611575178962655), Correlation_ID(18) +Record_ID(44), Domain(HSA_API_DOMAIN), Function(hsa_iterate_agents), Start_Timestamp(611575178953407), End_Timestamp(611575178962865), Correlation_ID(7) +Record_ID(46), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_allocate), Start_Timestamp(611575178963707), End_Timestamp(611575179001921), Correlation_ID(23) +Record_ID(48), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_allocate), Start_Timestamp(611575179002342), End_Timestamp(611575179019475), Correlation_ID(24) +Record_ID(50), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_allocate), Start_Timestamp(611575179019785), End_Timestamp(611575179038451), Correlation_ID(25) +Record_ID(52), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Start_Timestamp(611575179038812), End_Timestamp(611575179039153), Correlation_ID(26) +Record_ID(54), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Start_Timestamp(611575179041257), End_Timestamp(611575179043671), Correlation_ID(27) +Record_ID(56), Domain(HSA_API_DOMAIN), Function(hsa_agent_get_info), Start_Timestamp(611575179044112), End_Timestamp(611575179044353), Correlation_ID(28) +Record_ID(58), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_fill), Start_Timestamp(611575179045936), End_Timestamp(611575179047850), Correlation_ID(29) +Record_ID(60), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_fill), Start_Timestamp(611575179048090), End_Timestamp(611575184162319), Correlation_ID(30) +Record_ID(62), Domain(HSA_API_DOMAIN), Function(hsa_amd_agents_allow_access), Start_Timestamp(611575184164133), End_Timestamp(611575184170194), Correlation_ID(31) +Record_ID(64), Domain(HSA_API_DOMAIN), Function(hsa_signal_create), Start_Timestamp(611575184172248), End_Timestamp(611575184173681), Correlation_ID(32) +Record_ID(66), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_async_copy), Start_Timestamp(611575184174633), End_Timestamp(611575185129880), Correlation_ID(33) +Record_ID(69), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_copy_engine_status), Start_Timestamp(611575185131774), End_Timestamp(611575186734403), Correlation_ID(34) +Record_ID(71), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_async_copy_on_engine), Start_Timestamp(611575186736046), End_Timestamp(611575186742388), Correlation_ID(35) +Record_ID(73), Domain(HSA_API_DOMAIN), Function(hsa_signal_wait_relaxed), Start_Timestamp(611575186743440), End_Timestamp(611575186745304), Correlation_ID(36) +Record_ID(75), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_fill), Start_Timestamp(611575186748340), End_Timestamp(611575186792886), Correlation_ID(37) +Record_ID(77), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_fill), Start_Timestamp(611575186793246), End_Timestamp(611575186804628), Correlation_ID(38) +Record_ID(79), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_fill), Start_Timestamp(611575186804939), End_Timestamp(611575186815590), Correlation_ID(39) +Record_ID(81), Domain(HSA_API_DOMAIN), Function(hsa_amd_agents_allow_access), Start_Timestamp(611575186815940), End_Timestamp(611575186818124), Correlation_ID(40) +Record_ID(83), Domain(HSA_API_DOMAIN), Function(hsa_signal_create), Start_Timestamp(611575186818495), End_Timestamp(611575186819146), Correlation_ID(41) +Record_ID(85), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_async_copy), Start_Timestamp(611575186819497), End_Timestamp(611575186823815), Correlation_ID(42) +Record_ID(87), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_copy_engine_status), Start_Timestamp(611575186824166), End_Timestamp(611575186824577), Correlation_ID(43) +Record_ID(89), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_async_copy_on_engine), Start_Timestamp(611575186824877), End_Timestamp(611575186828645), Correlation_ID(44) +Record_ID(92), Domain(HSA_API_DOMAIN), Function(hsa_signal_wait_relaxed), Start_Timestamp(611575186828975), End_Timestamp(611575186837953), Correlation_ID(45) +Record_ID(94), Domain(HSA_API_DOMAIN), Function(hsa_signal_store_screlease), Start_Timestamp(611575186838834), End_Timestamp(611575186839025), Correlation_ID(46) +Record_ID(96), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_async_copy), Start_Timestamp(611575186839215), End_Timestamp(611575186842482), Correlation_ID(47) +Record_ID(99), Domain(HSA_API_DOMAIN), Function(hsa_signal_wait_relaxed), Start_Timestamp(611575186842832), End_Timestamp(611575186844946), Correlation_ID(48) +Record_ID(102), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_fill), Start_Timestamp(611575186846349), End_Timestamp(611575186862961), Correlation_ID(49) +Record_ID(104), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_fill), Start_Timestamp(611575186863242), End_Timestamp(611575191831478), Correlation_ID(50) +Record_ID(106), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_fill), Start_Timestamp(611575191832440), End_Timestamp(611575191851136), Correlation_ID(51) +Record_ID(108), Domain(HSA_API_DOMAIN), Function(hsa_amd_agents_allow_access), Start_Timestamp(611575191851537), End_Timestamp(611575191856146), Correlation_ID(52) +Record_ID(110), Domain(HSA_API_DOMAIN), Function(hsa_signal_create), Start_Timestamp(611575191856577), End_Timestamp(611575191857238), Correlation_ID(53) +Record_ID(112), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_async_copy), Start_Timestamp(611575191858962), End_Timestamp(611575191862969), Correlation_ID(54) +Record_ID(114), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_copy_engine_status), Start_Timestamp(611575191863330), End_Timestamp(611575191863961), Correlation_ID(55) +Record_ID(116), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_async_copy_on_engine), Start_Timestamp(611575191864292), End_Timestamp(611575191866526), Correlation_ID(56) +Record_ID(119), Domain(HSA_API_DOMAIN), Function(hsa_signal_wait_relaxed), Start_Timestamp(611575191866887), End_Timestamp(611575192213738), Correlation_ID(57) +Record_ID(121), Domain(HSA_API_DOMAIN), Function(hsa_amd_agents_allow_access), Start_Timestamp(611575192214118), End_Timestamp(611575192227644), Correlation_ID(58) +Record_ID(123), Domain(HSA_API_DOMAIN), Function(hsa_signal_store_screlease), Start_Timestamp(611575192228005), End_Timestamp(611575192228226), Correlation_ID(59) +Record_ID(125), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_async_copy), Start_Timestamp(611575192228506), End_Timestamp(611575193167662), Correlation_ID(60) +Record_ID(128), Domain(HSA_API_DOMAIN), Function(hsa_signal_wait_relaxed), Start_Timestamp(611575193168654), End_Timestamp(611575193243980), Correlation_ID(61) +Record_ID(130), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_free), Start_Timestamp(611575193246244), End_Timestamp(611575193271343), Correlation_ID(62) +Record_ID(132), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_free), Start_Timestamp(611575193272976), End_Timestamp(611575193274038), Correlation_ID(63) +Record_ID(134), Domain(HSA_API_DOMAIN), Function(hsa_amd_memory_pool_free), Start_Timestamp(611575193274318), End_Timestamp(611575193274890), Correlation_ID(64) diff --git a/tests-v2/featuretests/tracer/tracer_gtest.cpp b/tests-v2/featuretests/tracer/tracer_gtest.cpp index 4e506645de..82c03ba86a 100644 --- a/tests-v2/featuretests/tracer/tracer_gtest.cpp +++ b/tests-v2/featuretests/tracer/tracer_gtest.cpp @@ -23,7 +23,7 @@ THE SOFTWARE. #include #include #include "tracer_gtest.h" -#include "../utils/test_utils.h" + /** * Sets application enviornment by seting HSA_TOOLS_LIB. @@ -72,40 +72,16 @@ void ApplicationParser::SetApplicationEnv(const char* app_name, const char* trac * and saves them in a vector. */ void ApplicationParser::GetKernelInfoForRunningApplication( - std::vector* kernel_info_output) { - KernelInfo kinfo; + std::vector* kernel_info_output) { + tracer_kernel_info_t kinfo; for (std::string line : output_lines) { - // if (std::regex_match(line, std::regex("(Record)(.*)"))) { - // Record id - size_t found = line.find("Record"); - if (found != std::string::npos) { - int spos = found; - int epos = line.find(")", spos); - int length = std::string("Record").length(); - std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); - - kinfo.record_id = sub; + // Skip all the lines until "_DOMAIN" is found + if (line.empty() || line.find("_DOMAIN") == std::string::npos) { + continue; // Skip to the next line if "Dispatch_ID" is found } - // Kernel-Name - found = line.find("Function"); - if (found != std::string::npos) { - int spos = found; - int epos = line.find(")", spos); - int length = std::string("Function").length(); - std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); - kinfo.function = sub; - } - - // corelation-ids - found = line.find("Correlation_ID"); - if (found != std::string::npos) { - int spos = found; - int epos = line.find(")", spos); - int length = std::string("Correlation_ID").length(); - std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); - kinfo.corelation_id = sub; - } + // Parse individual values and store them in the dispatch struct + tokenize_tracer_output(line, kinfo); if (kinfo.record_id != "") { kernel_info_output->push_back(kinfo); @@ -117,8 +93,9 @@ void ApplicationParser::GetKernelInfoForRunningApplication( * Parses kernel-names from a pre-saved golden out files * and saves them in a vector. */ -void ApplicationParser::GetKernelInfoForGoldenOutput(const char* app_name, std::string file_name, - std::vector* kernel_info_output) { +void ApplicationParser::GetKernelInfoForGoldenOutput( + const char* app_name, std::string file_name, + std::vector* kernel_info_output) { std::string entry; std::string path = GetRunningPath("runTracerFeatureTests"); entry = path.append("apps/goldentraces/") + file_name; @@ -159,51 +136,22 @@ void ApplicationParser::ProcessApplication(std::stringstream& ss) { * Parses kernel-info for golden output file * and saves them in a vector. */ -void ApplicationParser::ParseKernelInfoFields(const std::string& s, - std::vector* kernel_info_output) { +void ApplicationParser::ParseKernelInfoFields( + const std::string& s, std::vector* kernel_info_output) { std::string line; - KernelInfo kinfo; + tracer_kernel_info_t kinfo; std::ifstream golden_file(s); while (!golden_file.eof()) { getline(golden_file, line); - // if (std::regex_match(line, std::regex("(Record)(.*)"))) { - // Record id - size_t found = line.find("Record"); - if (found != std::string::npos) { - int spos = found; - int epos = line.find(")", spos); - int length = std::string("Record").length(); - std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); - - kinfo.record_id = sub; - // kernel_info_output->push_back(kinfo); + // Skip all the lines until "_DOMAIN" is found + if (line.empty() || line.find("_DOMAIN") == std::string::npos) { + continue; // Skip to the next line if "Dispatch_ID" is found } - // Kernel-Name - found = line.find("Function"); - if (found != std::string::npos) { - int spos = found; - int epos = line.find(")", spos); - int length = std::string("Function").length(); - std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); + // Parse individual values and store them in the dispatch struct + tokenize_tracer_output(line, kinfo); - kinfo.function = sub; - // kernel_info_output->push_back(kinfo); - } - - // corealtion-ids - found = line.find("Correlation_ID"); - if (found != std::string::npos) { - int spos = found; - int epos = line.find(")", spos); - int length = std::string("Correlation_ID").length(); - std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); - - kinfo.corelation_id = sub; - // kernel_info_output->push_back(kinfo); - } - //} if (kinfo.record_id != "") { kernel_info_output->push_back(kinfo); } @@ -220,7 +168,7 @@ constexpr auto kGoldenOutputHelloworld = "hip_helloworld_golden_traces.txt"; class HelloWorldTest : public Tracertest { protected: - std::vector golden_kernel_info; + std::vector golden_kernel_info; void SetUp() { Tracertest::SetUp("tracer_hip_helloworld", "--hip-api "); GetKernelInfoForGoldenOutput("tracer_hip_helloworld", kGoldenOutputHelloworld, @@ -233,7 +181,7 @@ class HelloWorldTest : public Tracertest { // tracer output TEST_F(HelloWorldTest, WhenRunningTracerWithAppThenKernelInfoMatchWithGoldenOutput) { // kernel info in current profler run - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); @@ -245,7 +193,7 @@ TEST_F(HelloWorldTest, WhenRunningTracerWithAppThenKernelInfoMatchWithGoldenOutp // tracer output TEST_F(HelloWorldTest, WhenRunningTracerWithAppThenFunctionNamessMatchWithGoldenOutput) { // kernel info in current tracer run - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); @@ -258,7 +206,7 @@ TEST_F(HelloWorldTest, WhenRunningTracerWithAppThenFunctionNamessMatchWithGolden // tracer output TEST_F(HelloWorldTest, WhenRunningTracerWithAppThenKernelDurationShouldBePositive) { // kernel info in current tracer run - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); @@ -283,7 +231,7 @@ class AsyncCopyTest : public Tracertest { // tracer output TEST_F(AsyncCopyTest, WhenRunningTracerWithAppThenAsyncCopyOutputIsgenerated) { // kernel info in current profler run - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); @@ -292,14 +240,14 @@ TEST_F(AsyncCopyTest, WhenRunningTracerWithAppThenAsyncCopyOutputIsgenerated) { // Test:2 Matches coelation Ids TEST_F(AsyncCopyTest, WhenRunningTracerWithAppThenAsyncCorelationCountIsCorrect) { // kernel info in current profler run - std::vector current_kernel_info; + std::vector current_kernel_info; GetKernelInfoForRunningApplication(¤t_kernel_info); ASSERT_TRUE(current_kernel_info.size()); std::vector> corelation_pair{}; for (const auto& itr : current_kernel_info) { - if (itr.function.find("async_copy_on_engine") != std::string::npos) { + if (itr.domain.find("HSA_OPS_DOMAIN") != std::string::npos) { corelation_pair.push_back({itr.record_id, itr.corelation_id}); break; // we just want first occurance to test } diff --git a/tests-v2/featuretests/tracer/tracer_gtest.h b/tests-v2/featuretests/tracer/tracer_gtest.h index 0548d5d6ed..ee33006c02 100644 --- a/tests-v2/featuretests/tracer/tracer_gtest.h +++ b/tests-v2/featuretests/tracer/tracer_gtest.h @@ -32,6 +32,7 @@ THE SOFTWARE. #include #include #include +#include "../utils/test_utils.h" /* --------------------------------------------------------------------------*/ /** @@ -50,15 +51,6 @@ class ApplicationParser : public ::testing::Test { } virtual void TearDown() {} //!< This can be appended for other kernel info fields; eg: Agent-Name etc. - struct KernelInfo { - std::string record_id; - std::string domain; - std::string begin_time; - std::string end_time; - std::string corelation_id; - std::string roctx_id; - std::string function; - }; //!< saves lines of tracer output std::vector output_lines; @@ -70,11 +62,11 @@ class ApplicationParser : public ::testing::Test { //!< Parses kernel-info from a pre-saved golden out files // and saves them in a vector. void GetKernelInfoForGoldenOutput(const char* app_name, std::string filename, - std::vector* kernel_info_output); + std::vector* kernel_info_output); //!< Parses kernel-info after running tracer against curent application // and saves them in a vector. - void GetKernelInfoForRunningApplication(std::vector* kernel_info_output); + void GetKernelInfoForRunningApplication(std::vector* kernel_info_output); private: //!< Runs a given appllication and saves tracer output. @@ -84,7 +76,8 @@ class ApplicationParser : public ::testing::Test { //!< Parses kernel info fields from given input // i.e: kernel_names, kernel_duration - void ParseKernelInfoFields(const std::string& s, std::vector* kernel_info_output); + void ParseKernelInfoFields(const std::string& s, + std::vector* kernel_info_output); }; /* --------------------------------------------------------------------------*/ diff --git a/tests-v2/featuretests/utils/test_utils.cpp b/tests-v2/featuretests/utils/test_utils.cpp index 61811a5b20..8a56b5b987 100644 --- a/tests-v2/featuretests/utils/test_utils.cpp +++ b/tests-v2/featuretests/utils/test_utils.cpp @@ -79,6 +79,74 @@ bool is_installed_path() { return false; } +// tokenize profiler output +void tokenize_profiler_output(std::string line, profiler_kernel_info_t& kinfo) { + std::stringstream tokenStream(line); + std::string token; + std::getline(tokenStream, token, ','); + kinfo.record_id = token; + std::getline(tokenStream, token, ','); + kinfo.gpu_id = token; + std::getline(tokenStream, token, ','); + kinfo.queue_id = token; + std::getline(tokenStream, token, ','); + kinfo.queue_index = token; + std::getline(tokenStream, token, ','); + kinfo.process_id = token; + std::getline(tokenStream, token, ','); + kinfo.thread_id = token; + std::getline(tokenStream, token, ','); + kinfo.grid_size = token; + std::getline(tokenStream, token, ','); + kinfo.workgroup_size = token; + std::getline(tokenStream, token, ','); + kinfo.lds = token; + std::getline(tokenStream, token, ','); + kinfo.scratch_size = token; + std::getline(tokenStream, token, ','); + kinfo.arch_vgpr = token; + std::getline(tokenStream, token, ','); + kinfo.accum_vgpr = token; + std::getline(tokenStream, token, ','); + kinfo.sgpr = token; + std::getline(tokenStream, token, ','); + kinfo.wave_size = token; + std::getline(tokenStream, token, ','); + kinfo.kernel_name = token; + std::getline(tokenStream, token, ','); + kinfo.begin_time = token; + std::getline(tokenStream, token, ','); + kinfo.end_time = token; + std::getline(tokenStream, token, ','); + kinfo.counter = token; +} + +// tokenize tracer output +void tokenize_tracer_output(std::string line, tracer_kernel_info_t& kinfo) { + std::stringstream tokenStream(line); + std::string token; + std::getline(tokenStream, token, ','); + kinfo.record_id = token; + std::getline(tokenStream, token, ','); + kinfo.domain = token; + std::getline(tokenStream, token, ','); + kinfo.function = token; + std::getline(tokenStream, token, ','); + kinfo.operation = token; + std::getline(tokenStream, token, ','); + kinfo.kernel_name = token; + std::getline(tokenStream, token, ','); + kinfo.begin_time = token; + std::getline(tokenStream, token, ','); + kinfo.end_time = token; + std::getline(tokenStream, token, ','); + kinfo.corelation_id = token; + std::getline(tokenStream, token, ','); + kinfo.roctx_id = token; + std::getline(tokenStream, token, ','); + kinfo.roxtx_msg = token; +} + } // namespace utility } // namespace tests } // namespace rocprofiler diff --git a/tests-v2/featuretests/utils/test_utils.h b/tests-v2/featuretests/utils/test_utils.h index 786212cea8..65f0521b8c 100644 --- a/tests-v2/featuretests/utils/test_utils.h +++ b/tests-v2/featuretests/utils/test_utils.h @@ -29,6 +29,7 @@ THE SOFTWARE. #include #include #include +#include #include #include #include @@ -37,14 +38,55 @@ namespace rocprofiler { namespace tests { namespace utility { +typedef struct { + std::string record_id; + std::string gpu_id; + std::string queue_id; + std::string queue_index; + std::string process_id; + std::string thread_id; + std::string grid_size; + std::string workgroup_size; + std::string lds; + std::string scratch_size; + std::string arch_vgpr; + std::string accum_vgpr; + std::string sgpr; + std::string wave_size; + std::string kernel_name; + std::string begin_time; + std::string end_time; + std::string counter; +} profiler_kernel_info_t; + +typedef struct { + std::string record_id; + std::string domain; + std::string function; + std::string operation; + std::string kernel_name; + std::string begin_time; + std::string end_time; + std::string corelation_id; + std::string roctx_id; + std::string roxtx_msg; +} tracer_kernel_info_t; + // Get current running path std::string GetRunningPath(std::string string_to_erase); // Get Number of cores in the system int GetNumberOfCores(); +// Check if running path is /opt/rocm or not bool is_installed_path(); +// tokenize profiler output +void tokenize_profiler_output(std::string line, profiler_kernel_info_t& kinfo); + +// tokenize tracer output +void tokenize_tracer_output(std::string line, tracer_kernel_info_t& kinfo); + } // namespace utility } // namespace tests } // namespace rocprofiler @@ -56,5 +98,9 @@ int main(int argc, char** argv); using rocprofiler::tests::utility::GetNumberOfCores; using rocprofiler::tests::utility::GetRunningPath; using rocprofiler::tests::utility::is_installed_path; +using rocprofiler::tests::utility::profiler_kernel_info_t; +using rocprofiler::tests::utility::tokenize_profiler_output; +using rocprofiler::tests::utility::tokenize_tracer_output; +using rocprofiler::tests::utility::tracer_kernel_info_t; #endif // TESTS_FEATURETESTS_PROFILER_UTILS_TEST_UTILS_H_