diff --git a/projects/rocprofiler-sdk/source/lib/python/rocpd/csv.py b/projects/rocprofiler-sdk/source/lib/python/rocpd/csv.py index 0c6a5d3d9d..44a76a5ae4 100644 --- a/projects/rocprofiler-sdk/source/lib/python/rocpd/csv.py +++ b/projects/rocprofiler-sdk/source/lib/python/rocpd/csv.py @@ -2,7 +2,7 @@ ############################################################################### # MIT License # -# Copyright (c) 2023 Advanced Micro Devices, Inc. +# Copyright (c) 2025 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 @@ -23,14 +23,394 @@ # THE SOFTWARE. ############################################################################### +import os +import re + from .importer import RocpdImportData +from .query import export_sqlite_query from .time_window import apply_time_window from . import output_config from . import libpyrocpd +def write_sql_query_to_csv( + connection: RocpdImportData, + config, + query, + filename="", + postfix="trace", +) -> None: + """Write the contents of a SQL query to a CSV file in the specified output path.""" + + query_not_empty = f""" + SELECT EXISTS ( + {query} + ) + """ + + # just return if the result is empty + if not connection.execute(query_not_empty).fetchone()[0]: + return + + # call query module to export to csv + file_prefix = config.output_file + "_" if config.output_file else "" + file_postfix = "_" + postfix if postfix else "" + export_path = os.path.join( + config.output_path, f"{file_prefix}{filename}{file_postfix}.csv" + ) + + kwargs = {"title_columns": True} + export_sqlite_query( + connection, query, export_format="csv", export_path=export_path, **kwargs + ) + + +def write_agent_info_csv(importData, config) -> None: + + # Define mapping of output column name to JSON key + json_keys = [ + "node_id", + "logical_node_id", + "cpu_cores_count", + "simd_count", + "cpu_core_id_base", + "simd_id_base", + "max_waves_per_simd", + "lds_size_in_kb", + "gds_size_in_kb", + "num_gws", + "wave_front_size", + "num_xcc", + "cu_count", + "array_count", + "num_shader_banks", + "simd_arrays_per_engine", + "cu_per_simd_array", + "simd_per_cu", + "max_slots_scratch_cu", + "gfx_target_version", + "vendor_id", + "device_id", + "location_id", + "domain", + "drm_render_minor", + "num_sdma_engines", + "num_sdma_xgmi_engines", + "num_sdma_queues_per_engine", + "num_cp_queues", + "max_engine_clk_ccompute", + "max_engine_clk_fcompute", + "sdma_fw_version.uCodeSDMA AS Sdma_Fw_Version", + "fw_version.uCode AS Fw_Version", + "cu_per_engine", + "max_waves_per_cu", + "workgroup_max_size", + "family_id", + "grid_max_size", + "local_mem_size", + "hive_id", + "gpu_id", + "workgroup_max_dim.x AS Workgroup_Max_Dim_X", + "workgroup_max_dim.y AS Workgroup_Max_Dim_Y", + "workgroup_max_dim.z AS Workgroup_Max_Dim_Z", + "grid_max_dim.x AS Grid_Max_Dim_X", + "grid_max_dim.y AS Grid_Max_Dim_Y", + "grid_max_dim.z AS Grid_Max_Dim_Z", + "vendor_name", + "product_name", + ] + + # Build SELECT clause for json_extract columns + select_json = [] + for column in json_keys: + match = re.match(r"(.+?)\s+AS\s+(.+)", column, re.IGNORECASE) + column_name, column_alias = ( + (match.group(1), match.group(2)) if match else (column, column) + ) + select_json.append(f"json_extract(extdata, '$.{column_name}') AS {column_alias}") + + capabilities = [ + "HotPluggable", + "HSAMMUPresent", + "SharedWithGraphics", + "QueueSizePowerOfTwo", + "QueueSize32bit", + "QueueIdleEvent", + "VALimit", + "WatchPointsSupported", + "WatchPointsTotalBits", + "DoorbellType", + "AQLQueueDoubleMap", + "DebugTrapSupported", + "WaveLaunchTrapOverrideSupported", + "WaveLaunchModeSupported", + "PreciseMemoryOperationsSupported", + "DEPRECATED_SRAM_EDCSupport", + "Mem_EDCSupport", + "RASEventNotify", + "ASICRevision", + "SRAM_EDCSupport", + "SVMAPISupported", + "CoherentHostAccess", + "DebugSupportedFirmware", + "PreciseALUOperationsSupported", + "PerQueueResetSupported", + ] + + # Build SELECT clause for Capability columns + select_capability = [] + for capability in capabilities: + select_capability.append( + f"json_extract(extdata, '$.capability.{capability}') AS Cap_{capability}" + ) + + # Add non-JSON columns + fixed_keys = [ + "guid", + "type AS Agent_Type", + "name", + "model_name", + ] + + # to keep the right order + select_clause = ( + fixed_keys[:1] + + select_json[:2] + + fixed_keys[1:2] + + select_json[2:33] + + select_capability + + select_json[33:47] + + fixed_keys[2:3] + + select_json[47:] + + fixed_keys[3:4] + ) + + select_clause = ",\n ".join(select_clause) + + query = f""" + SELECT + {select_clause} + FROM "rocpd_info_agent" + """ + + write_sql_query_to_csv(importData, config, query, "agent_info", "") + + +def build_agent_id_string(agent_index_value, prefix=""): + + agent_prefix = prefix + "_" if prefix else "" + + if agent_index_value == libpyrocpd.agent_indexing.node: # absolute + return f"'Agent ' || {agent_prefix}agent_abs_index" + elif ( + agent_index_value == libpyrocpd.agent_indexing.logical_node + ): # relative (default) + return f"'Agent ' || {agent_prefix}agent_log_index" + elif ( + agent_index_value == libpyrocpd.agent_indexing.logical_node_type + ): # type-relative + return f"{agent_prefix}agent_type || ' ' || {agent_prefix}agent_type_index" + else: + return "" + + +def write_kernel_csv(importData, config) -> None: + + agent_id = build_agent_id_string(config.agent_index_value) + + if config.kernel_rename: + kernel_name = "region" + else: + kernel_name = "name" + + select_columns = [ + "guid", + "'KERNEL_DISPATCH' AS Kind", + f"{agent_id} AS Agent_Id", + "queue_id", + "stream_id", + "tid AS Thread_Id", + "dispatch_id", + "kernel_Id", + f"{kernel_name} AS Kernel_Name", + "stack_id AS Correlation_Id", + "start AS Start_Timestamp", + "end AS End_Timestamp", + "lds_size AS Lds_Block_Size", + "scratch_size", + "vgpr_count", + "accum_vgpr_count", + "sgpr_count", + "workgroup_x AS Workgroup_Size_X", + "workgroup_y AS Workgroup_Size_Y", + "workgroup_z AS Workgroup_Size_Z", + "grid_x AS Grid_Size_X", + "grid_y AS Grid_Size_Y", + "grid_z AS Grid_Size_Z", + ] + + aliased_headers = [] + for column in select_columns: + aliased_headers.append(column) + + select_clause = ",\n".join(aliased_headers) + + query = f""" + SELECT + {select_clause} + FROM "kernels" + ORDER BY + guid ASC, start ASC, end DESC + """ + write_sql_query_to_csv(importData, config, query, "kernel") + + +def write_memory_copy_csv(importData, config) -> None: + + src_agent_id = build_agent_id_string(config.agent_index_value, "src") + dst_agent_id = build_agent_id_string(config.agent_index_value, "dst") + + query = f""" + SELECT + guid, + 'MEMORY_COPY' AS Kind, + name AS Direction, + stream_id, + {src_agent_id} AS Source_Agent_Id, + {dst_agent_id} AS Destination_Agent_Id, + stack_id AS Correlation_Id, + start AS Start_Timestamp, + end AS End_Timestamp + FROM "memory_copies" + ORDER BY + guid ASC, start ASC, end DESC + """ + write_sql_query_to_csv(importData, config, query, "memory_copy") + + +def write_memory_allocation_csv(importData, config) -> None: + + agent_id = build_agent_id_string(config.agent_index_value) + + query = f""" + SELECT + guid, + 'MEMORY_ALLOCATION' AS Kind, + CASE + WHEN type = 'ALLOC' + THEN 'MEMORY_ALLOCATION_ALLOCATE' + ELSE 'MEMORY_ALLOCATION_' || type + END AS Operation, + CASE + WHEN type != 'FREE' + THEN {agent_id} + ELSE '"' + END AS Agent_Id, + size AS Allocation_Size, + '0x' || printf('%016X', address) AS Address, + stack_id AS Correlation_Id, + start AS Start_Timestamp, + end AS End_Timestamp + FROM "memory_allocations" + ORDER BY + guid ASC, start ASC, end DESC + """ + write_sql_query_to_csv(importData, config, query, "memory_allocation") + + +def write_counters_csv(importData, config) -> None: + + agent_id = build_agent_id_string(config.agent_index_value) + + select_columns = [ + "guid", + "stack_id AS Correlation_Id", + "dispatch_id", + f"{agent_id} AS Agent_Id", + "queue_id", + "pid AS Process_Id", + "tid AS Thread_Id", + "grid_size", + "kernel_id", + "kernel_name", + "workgroup_size", + "lds_block_size AS Lds_Block_Size", + "scratch_size", + "vgpr_count", + "accum_vgpr_count", + "sgpr_count", + "counter_name", + "value AS Counter_Value", + "start AS Start_Timestamp", + "end AS End_Timestamp", + ] + + aliased_headers = [] + for column in select_columns: + aliased_headers.append(column) + + select_clause = ",\n".join(aliased_headers) + + query = f""" + SELECT + {select_clause} + FROM "counters_collection" + ORDER BY + guid ASC, start ASC, end DESC + """ + write_sql_query_to_csv(importData, config, query, "counter_collection") + + +def write_scratch_memory_csv(importData, config) -> None: + + agent_id = build_agent_id_string(config.agent_index_value) + + query = f""" + SELECT + guid, + 'SCRATCH_MEMORY' AS Kind, + 'SCRATCH_MEMORY_' || operation AS Operation, + {agent_id} AS Agent_Id, + queue_id, + tid AS Thread_Id, + alloc_flags, + start AS Start_Timestamp, + end AS End_Timestamp + FROM "scratch_memory" + ORDER BY + guid ASC, start ASC, end DESC + """ + write_sql_query_to_csv(importData, config, query, "scratch_memory") + + +def write_region_csv(importData, config) -> None: + + query = """ + SELECT + guid, + category AS Domain, + name AS Function, + pid AS Process_Id, + tid AS Thread_Id, + stack_id AS Correlation_Id, + start AS Start_Timestamp, + end AS End_Timestamp + FROM "regions" + ORDER BY + guid ASC, start ASC, end DESC + """ + write_sql_query_to_csv(importData, config, query, "regions") + + def write_csv(importData, config): - return libpyrocpd.write_csv(importData, config) + + write_agent_info_csv(importData, config) + write_counters_csv(importData, config) + write_kernel_csv(importData, config) + write_memory_allocation_csv(importData, config) + write_memory_copy_csv(importData, config) + write_region_csv(importData, config) + write_scratch_memory_csv(importData, config) def execute(input, config=None, window_args=None, **kwargs): diff --git a/projects/rocprofiler-sdk/source/lib/python/rocpd/libpyrocpd.cpp b/projects/rocprofiler-sdk/source/lib/python/rocpd/libpyrocpd.cpp index 0f62325724..02e8e60691 100644 --- a/projects/rocprofiler-sdk/source/lib/python/rocpd/libpyrocpd.cpp +++ b/projects/rocprofiler-sdk/source/lib/python/rocpd/libpyrocpd.cpp @@ -23,7 +23,6 @@ #include "libpyrocpd.hpp" #include "lib/output/format_path.hpp" #include "lib/python/rocpd/source/common.hpp" -#include "lib/python/rocpd/source/csv.hpp" #include "lib/python/rocpd/source/functions.hpp" #include "lib/python/rocpd/source/interop.hpp" #include "lib/python/rocpd/source/otf2.hpp" @@ -505,112 +504,6 @@ PYBIND11_MODULE(libpyrocpd, pyrocpd) }, "Write pftrace output file from rocpd SQLite3 database"); - pyrocpd.def( - "write_csv", - [](rocpd::RocpdImportData& data, const rocprofiler::tool::output_config& output_cfg) { - auto sqlgen_csv = common::simple_timer{ - fmt::format("CSV generation from {} SQL database(s)", data.size())}; - - if(data.empty()) return; - - auto csv_manager = rocpd::output::CsvManager{output_cfg}; - - for(auto obj : {data.connection}) - { - auto* conn = rocpd::interop::get_connection(std::move(obj)); - auto nodes = rocpd::read(conn); - - for(const auto& nitr : nodes) - { - auto agents = rocpd::read( - conn, fmt::format("WHERE guid = '{}' AND nid = {}", nitr.guid, nitr.id)); - auto processes = rocpd::read( - conn, fmt::format("WHERE guid = '{}' AND nid = {}", nitr.guid, nitr.id)); - - for(const auto& pitr : processes) - { - ROCP_FATAL_IF(pitr.nid != nitr.id || pitr.guid != nitr.guid) - << fmt::format("Found process with a mismatched nid/guid. process: " - "{}/{} vs. node: {}/{}", - pitr.nid, - pitr.guid, - nitr.id, - nitr.guid); - auto _sqlgen_csv = common::simple_timer{fmt::format( - "CSV generation from SQL for process {} (total)", pitr.pid)}; - - auto select_guid_nid_pid = [&nitr, &pitr](std::string_view tbl, - std::string_view - where_extra_condition = {}) { - return fmt::format( - "SELECT * FROM {} WHERE guid = '{}' AND nid = {} AND pid = {} {}", - tbl, - pitr.guid, - nitr.id, - pitr.pid, - where_extra_condition); - }; - - rocpd::output::write_agent_info_csv(csv_manager, agents); - - constexpr auto region_order_by = "start ASC, end DESC"; - - auto kernels = rocpd::sql_generator{ - conn, select_guid_nid_pid("kernels"), region_order_by}; - auto memory_copies = rocpd::sql_generator{ - conn, select_guid_nid_pid("memory_copies"), region_order_by}; - auto memory_allocations = - rocpd::sql_generator{ - conn, select_guid_nid_pid("memory_allocations"), region_order_by}; - auto hip_api_calls = rocpd::sql_generator{ - conn, - select_guid_nid_pid("regions", "AND category LIKE 'HIP_%'"), - region_order_by}; - auto hsa_api_calls = rocpd::sql_generator{ - conn, - select_guid_nid_pid("regions", "AND category LIKE 'HSA_%'"), - region_order_by}; - auto marker_api_calls = rocpd::sql_generator{ - conn, - select_guid_nid_pid("regions_and_samples", - "AND category LIKE 'MARKER_%'"), - region_order_by}; - auto counters_calls = rocpd::sql_generator{ - conn, select_guid_nid_pid("counters_collection"), region_order_by}; - auto scratch_memory_calls = - rocpd::sql_generator{ - conn, select_guid_nid_pid("scratch_memory"), region_order_by}; - auto rccl_calls = rocpd::sql_generator{ - conn, - select_guid_nid_pid("regions", "AND category LIKE 'RCCL_%'"), - region_order_by}; - auto rocdecode_calls = rocpd::sql_generator{ - conn, - select_guid_nid_pid("regions", "AND category LIKE 'ROCDECODE_%'"), - region_order_by}; - auto rocjpeg_calls = rocpd::sql_generator{ - conn, - select_guid_nid_pid("regions", "AND category LIKE 'ROCJPEG_%'"), - region_order_by}; - - rocpd::output::write_csvs(csv_manager, - kernels, - memory_copies, - memory_allocations, - hip_api_calls, - hsa_api_calls, - marker_api_calls, - counters_calls, - scratch_memory_calls, - rccl_calls, - rocdecode_calls, - rocjpeg_calls); - } - } - } - }, - "Write trace data to CSV files"); - pyrocpd.def( "write_otf2", [](rocpd::RocpdImportData& data, const tool::output_config& output_cfg) { diff --git a/projects/rocprofiler-sdk/source/lib/python/rocpd/query.py b/projects/rocprofiler-sdk/source/lib/python/rocpd/query.py index ab369326db..c4099f2994 100644 --- a/projects/rocprofiler-sdk/source/lib/python/rocpd/query.py +++ b/projects/rocprofiler-sdk/source/lib/python/rocpd/query.py @@ -42,6 +42,7 @@ def export_sqlite_query( export_format: Optional[str] = None, export_path: Optional[str] = None, dashboard_template_path: Optional[str] = None, + **kwargs: Optional[dict], ) -> Optional[str]: """ Execute a SQLite query and print it to console. @@ -99,7 +100,21 @@ def export_sqlite_query( # 3) Export based on format if export_format == "csv": - df.to_csv(export_path, index=False) + import csv + + cols = [f"{itr}" for itr in df.columns.tolist()] + col_names = ( + [f"{itr}".title() for itr in cols] + if kwargs.get("title_columns", True) + else cols[:] + ) + df.to_csv( + export_path, + index=False, + columns=cols, + header=col_names, + quoting=csv.QUOTE_NONNUMERIC, + ) elif export_format == "html": write_export(df.to_html(index=False)) diff --git a/projects/rocprofiler-sdk/source/lib/python/rocpd/source/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/python/rocpd/source/CMakeLists.txt index 99a6927b43..70ac637c1f 100644 --- a/projects/rocprofiler-sdk/source/lib/python/rocpd/source/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/python/rocpd/source/CMakeLists.txt @@ -2,19 +2,10 @@ # libpyrocpd python binding sources # -set(libpyrocpd_source_headers - common.hpp - functions.hpp - interop.hpp - perfetto.hpp - csv.hpp - otf2.hpp - sql_generator.hpp - pysqlite_Connection.h - types.hpp) +set(libpyrocpd_source_headers common.hpp functions.hpp interop.hpp perfetto.hpp otf2.hpp + sql_generator.hpp pysqlite_Connection.h types.hpp) -set(libpyrocpd_source_sources csv.cpp functions.cpp interop.cpp otf2.cpp perfetto.cpp - types.cpp) +set(libpyrocpd_source_sources functions.cpp interop.cpp otf2.cpp perfetto.cpp types.cpp) foreach(_PYTHON_VERSION ${ROCPROFILER_PYTHON_VERSIONS}) rocprofiler_rocpd_python_bindings_target_sources( diff --git a/projects/rocprofiler-sdk/source/lib/python/rocpd/source/csv.cpp b/projects/rocprofiler-sdk/source/lib/python/rocpd/source/csv.cpp deleted file mode 100644 index 57bb7232be..0000000000 --- a/projects/rocprofiler-sdk/source/lib/python/rocpd/source/csv.cpp +++ /dev/null @@ -1,743 +0,0 @@ -// MIT License -// -// Copyright (c) 2025 Advanced Micro Devices, Inc. All Rights Reserved. -// -// Permission is hereby granted, free of charge, to any person obtaining a copy -// of this software and associated documentation files (the "Software"), to deal -// in the Software without restriction, including without limitation the rights -// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -// copies of the Software, and to permit persons to whom the Software is -// furnished to do so, subject to the following conditions: -// -// The above copyright notice and this permission notice shall be included in -// all copies or substantial portions of the Software. -// -// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -// SOFTWARE. - -#include "lib/python/rocpd/source/csv.hpp" - -#include "lib/common/defines.hpp" -#include "lib/common/hasher.hpp" -#include "lib/common/mpl.hpp" -#include "lib/output/csv.hpp" -#include "lib/output/csv_output_file.hpp" -#include "lib/output/generator.hpp" -#include "lib/output/metadata.hpp" -#include "lib/output/node_info.hpp" -#include "lib/output/output_config.hpp" -#include "lib/output/output_stream.hpp" -#include "lib/output/sql/common.hpp" -#include "lib/output/stream_info.hpp" -#include "lib/rocprofiler-sdk-tool/config.hpp" - -#include - -#include - -#include -#include -#include -#include -#include -#include -#include - -namespace fs = std::filesystem; - -namespace -{ -const std::string STATS_HEADER = "\"Name\",\"Calls\",\"TotalDurationNs\"," - "\"AverageNs\",\"Percentage\",\"MinNs\",\"MaxNs\",\"StdDev\""; -const std::string API_TRACE_HEADER = - "\"Guid\",\"Domain\",\"Function\",\"Process_Id\",\"Thread_Id\"," - "\"Correlation_Id\",\"Start_Timestamp\",\"End_Timestamp\""; -} // namespace - -namespace rocpd -{ -namespace output -{ -CsvManager::CsvManager(rocprofiler::tool::output_config output_cfg) -: config{std::move(output_cfg)} -{ - if(!ensure_output_directory()) - { - ROCP_ERROR << "Failed to create csv output directory: " << config.output_path; - return; - } - - this->csv_configs = { - {CsvType::KERNEL_DISPATCH, - {"kernel_trace.csv", - "\"Guid\",\"Kind\",\"Agent_Id\",\"Queue_Id\",\"Stream_Id\",\"Thread_Id\",\"Dispatch_Id\"," - "\"Kernel_Id\",\"Kernel_Name\",\"Correlation_Id\",\"Start_Timestamp\",\"End_Timestamp\"," - "\"LDS_Block_Size\",\"Scratch_Size\",\"VGPR_Count\",\"Accum_VGPR_Count\",\"SGPR_Count\"," - "\"Workgroup_Size_X\",\"Workgroup_Size_Y\",\"Workgroup_Size_Z\"," - "\"Grid_Size_X\",\"Grid_Size_Y\",\"Grid_Size_Z\""}}, - {CsvType::MEMORY_COPY, - {"memory_copy_trace.csv", - "\"Guid\",\"Kind\",\"Direction\",\"Stream_Id\",\"Source_Agent_Id\"," - "\"Destination_Agent_" - "Id\"," - "\"Correlation_Id\",\"Start_Timestamp\",\"End_Timestamp\""}}, - {CsvType::MEMORY_ALLOCATION, - {"memory_allocation_trace.csv", - "\"Guid\",\"Kind\",\"Operation\",\"Agent_Id\",\"Allocation_Size\"," - "\"Address\"," - "\"Correlation_Id\",\"Start_Timestamp\",\"End_Timestamp\""}}, - {CsvType::SCRATCH_MEMORY, - {"scratch_memory_trace.csv", - "\"Kind\",\"Operation\",\"Agent_Id\",\"Queue_Id\",\"Thread_Id\"," - "\"Alloc_Flags\",\"Start_" - "Timestamp\",\"End_Timestamp\""}}, - - {CsvType::HIP_API, {"hip_api_trace.csv", API_TRACE_HEADER}}, - {CsvType::HSA_CSV_API, {"hsa_api_trace.csv", API_TRACE_HEADER}}, - {CsvType::MARKER, {"marker_api_trace.csv", API_TRACE_HEADER}}, - {CsvType::RCCL_API, {"rccl_api_trace.csv", API_TRACE_HEADER}}, - {CsvType::ROCDECODE_API, {"rocdecode_api_trace.csv", API_TRACE_HEADER}}, - {CsvType::ROCJPEG_API, {"rocjpeg_api_trace.csv", API_TRACE_HEADER}}, - - {CsvType::COUNTER, - {"counter_collection.csv", - "\"Pid\",\"Correlation_Id\",\"Dispatch_Id\",\"Agent_Id\",\"Queue_Id\"," - "\"Process_Id\"," - "\"Thread_Id\"," - "\"Grid_Size\",\"Kernel_Id\",\"Kernel_Name\",\"Workgroup_Size\",\"LDS_" - "Block_Size\"," - "\"Scratch_Size\",\"VGPR_Count\",\"Accum_VGPR_Count\",\"SGPR_Count\"," - "\"Counter_Name\",\"Counter_Value\",\"Start_Timestamp\",\"End_" - "Timestamp\""}}, - }; -} - -bool -CsvManager::ensure_output_directory() const -{ - try - { - fs::create_directories(config.output_path); - return true; - } catch(const std::exception& e) - { - ROCP_ERROR << "Failed to create directory: " << e.what(); - return false; - } -} - -CsvManager::~CsvManager() -{ - for(auto& [type, stream] : streams) - { - if(stream.is_open()) - { - stream.flush(); - stream.close(); - } - } -} - -std::ofstream& -CsvManager::get_stream(CsvType type) -{ - return streams[type]; -} - -bool -CsvManager::has_stream(CsvType type) const -{ - return streams.count(type) != 0u && streams.at(type).is_open(); -} - -bool -CsvManager::initialize_csv_file(CsvType type) -{ - if(has_stream(type)) return true; - - if(csv_configs.find(type) == csv_configs.end()) - { - ROCP_ERROR << "No CSV configuration found for type: " << static_cast(type); - return false; - } - - const auto& cfg = csv_configs[type]; - - fs::path output_dir = config.output_path; - fs::path filename = - config.output_file.empty() ? cfg.filename : config.output_file + "_" + cfg.filename; - - file_paths[type] = (output_dir / filename).string(); - - auto& path = file_paths[type]; - auto& stream = streams[type]; - - stream.open(path, std::ios::out); - if(!stream.is_open()) - { - ROCP_ERROR << "Failed to open CSV output file: " << path; - return false; - } - - stream << cfg.header << '\n'; - return true; -} - -template -bool -has_any_data(const rocprofiler::tool::generator& data_gen) -{ - for(auto ditr : data_gen) - { - auto gen = data_gen.get(ditr); - if(begin(gen) != end(gen)) - { - return true; - } - } - return false; -} - -template -void -process_data_to_csv(CsvManager& csv_manager, - CsvType csv_type, - const rocprofiler::tool::generator& data_gen, - Processor process_func) -{ - if(!has_any_data(data_gen)) return; - - if(!csv_manager.initialize_csv_file(csv_type)) return; - - for(auto ditr : data_gen) - { - auto gen = data_gen.get(ditr); - for(auto it = begin(gen); it != end(gen); ++it) - { - process_func(csv_manager, csv_type, *it); - } - } -} - -void -write_kernel_csv( - CsvManager& csv_manager, - const rocprofiler::tool::generator& kernel_dispatch_gen) -{ - process_data_to_csv( - csv_manager, - CsvType::KERNEL_DISPATCH, - kernel_dispatch_gen, - [](CsvManager& cm, CsvType type, const rocpd::types::kernel_dispatch& kernel) { - std::string kernel_identifier = - (cm.config.kernel_rename && !kernel.region.empty()) ? kernel.region : kernel.name; - std::string agent_identifier = create_agent_index(cm.config.agent_index_value, - kernel.agent_abs_index, - kernel.agent_log_index, - kernel.agent_type_index, - std::string_view(kernel.agent_type)) - .as_string(); - - cm.write_line(type, - fmt::format("\"{}\"", kernel.guid), - fmt::format("\"{}\"", "KERNEL_DISPATCH"), - fmt::format("\"{}\"", agent_identifier), - kernel.queue_id, - kernel.stream_id, - kernel.tid, - kernel.dispatch_id, - kernel.kernel_id, - fmt::format("\"{}\"", kernel_identifier), - kernel.stack_id, - kernel.start, - kernel.end, - kernel.lds_size, - kernel.scratch_size, - kernel.vgpr_count, - kernel.accum_vgpr_count, - kernel.sgpr_count, - kernel.workgroup_size.x, - kernel.workgroup_size.y, - kernel.workgroup_size.z, - kernel.grid_size.x, - kernel.grid_size.y, - kernel.grid_size.z); - }); -} - -void -write_memory_copy_csv( - CsvManager& csv_manager, - const rocprofiler::tool::generator& memory_copies_gen) -{ - process_data_to_csv(csv_manager, - CsvType::MEMORY_COPY, - memory_copies_gen, - [](CsvManager& cm, CsvType type, const rocpd::types::memory_copies& mcopy) { - std::string src_agent_identifier = - create_agent_index(cm.config.agent_index_value, - mcopy.src_agent_abs_index, - mcopy.src_agent_log_index, - mcopy.src_agent_type_index, - std::string_view(mcopy.src_agent_type)) - .as_string(); - - std::string dst_agent_identifier = - create_agent_index(cm.config.agent_index_value, - mcopy.dst_agent_abs_index, - mcopy.dst_agent_log_index, - mcopy.dst_agent_type_index, - std::string_view(mcopy.dst_agent_type)) - .as_string(); - - cm.write_line(type, - fmt::format("\"{}\"", mcopy.guid), - fmt::format("\"{}\"", "MEMORY_COPY"), - fmt::format("\"{}\"", mcopy.name), - mcopy.stream_id, - fmt::format("\"{}\"", src_agent_identifier), - fmt::format("\"{}\"", dst_agent_identifier), - mcopy.stack_id, - mcopy.start, - mcopy.end); - }); -} - -void -write_memory_allocation_csv( - CsvManager& csv_manager, - const rocprofiler::tool::generator& memory_alloc_gen) -{ - process_data_to_csv( - csv_manager, - CsvType::MEMORY_ALLOCATION, - memory_alloc_gen, - [](CsvManager& cm, CsvType type, const rocpd::types::memory_allocation& malloc) { - std::string normalized_type = malloc.type; - if(normalized_type == "ALLOC") - { - normalized_type = "ALLOCATE"; - } - - std::string operation = fmt::format("MEMORY_ALLOCATION_{}", normalized_type); - - std::string agent_identifier = create_agent_index(cm.config.agent_index_value, - malloc.agent_abs_index, - malloc.agent_log_index, - malloc.agent_type_index, - std::string_view(malloc.agent_type)) - .as_string(); - - std::string agent_id = - operation != "MEMORY_ALLOCATION_FREE" ? agent_identifier : "\"\""; - std::string address = fmt::format("\"0x{:016x}\"", malloc.address); - - cm.write_line(type, - fmt::format("\"{}\"", malloc.guid), - fmt::format("\"{}\"", "MEMORY_ALLOCATION"), - fmt::format("\"{}\"", operation), - fmt::format("\"{}\"", agent_id), - malloc.size, - address, - malloc.stack_id, - malloc.start, - malloc.end); - }); -} - -void -write_scratch_memory_csv( - CsvManager& csv_manager, - const rocprofiler::tool::generator& scratch_memory_gen) -{ - process_data_to_csv( - csv_manager, - CsvType::SCRATCH_MEMORY, - scratch_memory_gen, - [](CsvManager& cm, CsvType type, const rocpd::types::scratch_memory& scratch_mem) { - std::string agent_identifier = - create_agent_index(cm.config.agent_index_value, - scratch_mem.agent_abs_index, - scratch_mem.agent_log_index, - scratch_mem.agent_type_index, - std::string_view(scratch_mem.agent_type)) - .as_string(); - - cm.write_line(type, - fmt::format("\"{}\"", "SCRATCH_MEMORY"), - fmt::format("\"SCRATCH_MEMORY_{}\"", scratch_mem.operation), - fmt::format("\"{}\"", agent_identifier), - scratch_mem.queue_id, - scratch_mem.tid, - scratch_mem.alloc_flags, - scratch_mem.start, - scratch_mem.end); - }); -} - -void -write_hip_api_csv(CsvManager& csv_manager, - const rocprofiler::tool::generator& hip_api_gen) -{ - process_data_to_csv(csv_manager, - CsvType::HIP_API, - hip_api_gen, - [](CsvManager& cm, CsvType type, const rocpd::types::region& api) { - // Skip entries that are not HIP API calls - if(api.category.find("HIP_") != 0) return; - - cm.write_line(type, - fmt::format("\"{}\"", api.guid), - fmt::format("\"{}\"", api.category), - fmt::format("\"{}\"", api.name), - api.pid, - api.tid, - api.stack_id, - api.start, - api.end); - }); -} - -void -write_hsa_api_csv(CsvManager& csv_manager, - const rocprofiler::tool::generator& hsa_api_gen) -{ - process_data_to_csv(csv_manager, - CsvType::HSA_CSV_API, - hsa_api_gen, - [](CsvManager& cm, CsvType type, const rocpd::types::region& api) { - // Skip entries that are not HSA API calls - if(api.category.find("HSA_") != 0) return; - - cm.write_line(type, - fmt::format("\"{}\"", api.guid), - fmt::format("\"{}\"", api.category), - fmt::format("\"{}\"", api.name), - api.pid, - api.tid, - api.stack_id, - api.start, - api.end); - }); -} - -void -write_marker_api_csv(CsvManager& csv_manager, - const rocprofiler::tool::generator& marker_api_gen) -{ - namespace tool = ::rocprofiler::tool; - - if(marker_api_gen.empty()) return; - - using marker_csv_encoder = tool::csv::csv_encoder<8>; - - auto ofs = tool::csv_output_file{csv_manager.config, - domain_type::MARKER, - marker_csv_encoder{}, - {"Guid", - "Domain", - "Function", - "Process_Id", - "Thread_Id", - "Correlation_Id", - "Start_Timestamp", - "End_Timestamp"}}; - - // write samples first and ignore the timestamp ordering w.r.t. regions for now - for(auto ditr : marker_api_gen) - { - for(const auto& record : marker_api_gen.get(ditr)) - { - auto row_ss = std::stringstream{}; - auto _name = record.name; - - if(record.has_extdata()) - { - if(auto _extdata = record.get_extdata(); !_extdata.message.empty()) - _name = _extdata.message; - } - - marker_csv_encoder::write_row(row_ss, - record.guid, - record.category, - _name, - record.pid, - record.tid, - record.stack_id, - record.start, - record.end); - - ofs << row_ss.str(); - } - } -} - -void -write_rccl_api_csv(CsvManager& csv_manager, - const rocprofiler::tool::generator& rccl_api_gen) -{ - process_data_to_csv(csv_manager, - CsvType::RCCL_API, - rccl_api_gen, - [](CsvManager& cm, CsvType type, const rocpd::types::region& api) { - if(api.category.find("RCCL_") != 0) return; - - cm.write_line(type, - fmt::format("\"{}\"", api.guid), - fmt::format("\"{}\"", api.category), - fmt::format("\"{}\"", api.name), - api.pid, - api.tid, - api.stack_id, - api.start, - api.end); - }); -} - -void -write_rocdecode_api_csv(CsvManager& csv_manager, - const rocprofiler::tool::generator& rocdecode_api_gen) -{ - process_data_to_csv(csv_manager, - CsvType::ROCDECODE_API, - rocdecode_api_gen, - [](CsvManager& cm, CsvType type, const rocpd::types::region& api) { - if(api.category.find("ROCDECODE_") != 0) return; - - cm.write_line(type, - fmt::format("\"{}\"", api.guid), - fmt::format("\"{}\"", api.category), - fmt::format("\"{}\"", api.name), - api.pid, - api.tid, - api.stack_id, - api.start, - api.end); - }); -} - -void -write_rocjpeg_api_csv(CsvManager& csv_manager, - const rocprofiler::tool::generator& rocjpeg_api_gen) -{ - process_data_to_csv(csv_manager, - CsvType::ROCJPEG_API, - rocjpeg_api_gen, - [](CsvManager& cm, CsvType type, const rocpd::types::region& api) { - if(api.category.find("ROCJPEG_") != 0) return; - - cm.write_line(type, - fmt::format("\"{}\"", api.guid), - fmt::format("\"{}\"", api.category), - fmt::format("\"{}\"", api.name), - api.pid, - api.tid, - api.stack_id, - api.start, - api.end); - }); -} - -void -write_agent_info_csv(CsvManager& csv_manager, const std::vector& agents) -{ - if(agents.empty()) return; - - namespace tool = ::rocprofiler::tool; - using agent_info_csv_encoder = tool::csv::csv_encoder<54>; - - auto ofs = tool::csv_output_file{csv_manager.config, - "agent_info", - agent_info_csv_encoder{}, - {"Guid", - "Node_Id", - "Logical_Node_Id", - "Agent_Type", - "Cpu_Cores_Count", - "Simd_Count", - "Cpu_Core_Id_Base", - "Simd_Id_Base", - "Max_Waves_Per_Simd", - "Lds_Size_In_Kb", - "Gds_Size_In_Kb", - "Num_Gws", - "Wave_Front_Size", - "Num_Xcc", - "Cu_Count", - "Array_Count", - "Num_Shader_Banks", - "Simd_Arrays_Per_Engine", - "Cu_Per_Simd_Array", - "Simd_Per_Cu", - "Max_Slots_Scratch_Cu", - "Gfx_Target_Version", - "Vendor_Id", - "Device_Id", - "Location_Id", - "Domain", - "Drm_Render_Minor", - "Num_Sdma_Engines", - "Num_Sdma_Xgmi_Engines", - "Num_Sdma_Queues_Per_Engine", - "Num_Cp_Queues", - "Max_Engine_Clk_Ccompute", - "Max_Engine_Clk_Fcompute", - "Sdma_Fw_Version", - "Fw_Version", - "Capability", - "Cu_Per_Engine", - "Max_Waves_Per_Cu", - "Family_Id", - "Workgroup_Max_Size", - "Grid_Max_Size", - "Local_Mem_Size", - "Hive_Id", - "Gpu_Id", - "Workgroup_Max_Dim_X", - "Workgroup_Max_Dim_Y", - "Workgroup_Max_Dim_Z", - "Grid_Max_Dim_X", - "Grid_Max_Dim_Y", - "Grid_Max_Dim_Z", - "Name", - "Vendor_Name", - "Product_Name", - "Model_Name"}}; - - for(const auto& itr : agents) - { - auto row_ss = std::stringstream{}; - agent_info_csv_encoder::write_row(row_ss, - itr.guid, - itr.node_id, - itr.logical_node_id, - itr.type, - itr.cpu_cores_count, - itr.simd_count, - itr.cpu_core_id_base, - itr.simd_id_base, - itr.max_waves_per_simd, - itr.lds_size_in_kb, - itr.gds_size_in_kb, - itr.num_gws, - itr.wave_front_size, - itr.num_xcc, - itr.cu_count, - itr.array_count, - itr.num_shader_banks, - itr.simd_arrays_per_engine, - itr.cu_per_simd_array, - itr.simd_per_cu, - itr.max_slots_scratch_cu, - itr.gfx_target_version, - itr.vendor_id, - itr.device_id, - itr.location_id, - itr.domain, - itr.drm_render_minor, - itr.num_sdma_engines, - itr.num_sdma_xgmi_engines, - itr.num_sdma_queues_per_engine, - itr.num_cp_queues, - itr.max_engine_clk_ccompute, - itr.max_engine_clk_fcompute, - itr.sdma_fw_version.Value, - itr.fw_version.Value, - itr.capability.Value, - itr.cu_per_engine, - itr.max_waves_per_cu, - itr.family_id, - itr.workgroup_max_size, - itr.grid_max_size, - itr.local_mem_size, - itr.hive_id, - itr.gpu_id, - itr.workgroup_max_dim.x, - itr.workgroup_max_dim.y, - itr.workgroup_max_dim.z, - itr.grid_max_dim.x, - itr.grid_max_dim.y, - itr.grid_max_dim.z, - itr.name, - itr.vendor_name, - itr.product_name, - itr.model_name); - ofs << row_ss.str(); - } -} - -void -write_counters_csv(CsvManager& csv_manager, - const rocprofiler::tool::generator& counter_gen) -{ - process_data_to_csv(csv_manager, - CsvType::COUNTER, - counter_gen, - [](CsvManager& cm, CsvType type, const rocpd::types::counter& counter) { - std::string agent_identifier = - create_agent_index(cm.config.agent_index_value, - counter.agent_abs_index, - counter.agent_log_index, - counter.agent_type_index, - std::string_view(counter.agent_type)) - .as_string(); - - cm.write_line(type, - counter.guid, - counter.stack_id, - counter.dispatch_id, - fmt::format("\"{}\"", agent_identifier), - counter.queue_id, - counter.pid, - counter.tid, - counter.grid_size, - counter.kernel_id, - fmt::format("\"{}\"", counter.kernel_name), - counter.workgroup_size, - counter.lds_block_size, - counter.scratch_size, - counter.vgpr_count, - counter.accum_vgpr_count, - counter.sgpr_count, - fmt::format("\"{}\"", counter.counter_name), - counter.value, - counter.start, - counter.end); - }); -} - -void -write_csvs(CsvManager& csv_manager, - const rocprofiler::tool::generator& kernel_dispatch, - const rocprofiler::tool::generator& memory_copies, - const rocprofiler::tool::generator& memory_allocations, - const rocprofiler::tool::generator& hip_api_calls, - const rocprofiler::tool::generator& hsa_api_calls, - const rocprofiler::tool::generator& marker_api_calls, - const rocprofiler::tool::generator& counters_calls, - const rocprofiler::tool::generator& scratch_memory_calls, - const rocprofiler::tool::generator& rccl_calls, - const rocprofiler::tool::generator& rocdecode_calls, - const rocprofiler::tool::generator& rocjpeg_calls) -{ - rocpd::output::write_kernel_csv(csv_manager, kernel_dispatch); - rocpd::output::write_memory_copy_csv(csv_manager, memory_copies); - rocpd::output::write_memory_allocation_csv(csv_manager, memory_allocations); - rocpd::output::write_hip_api_csv(csv_manager, hip_api_calls); - rocpd::output::write_hsa_api_csv(csv_manager, hsa_api_calls); - rocpd::output::write_marker_api_csv(csv_manager, marker_api_calls); - - rocpd::output::write_counters_csv(csv_manager, counters_calls); - rocpd::output::write_scratch_memory_csv(csv_manager, scratch_memory_calls); - rocpd::output::write_rccl_api_csv(csv_manager, rccl_calls); - - rocpd::output::write_rocdecode_api_csv(csv_manager, rocdecode_calls); - rocpd::output::write_rocjpeg_api_csv(csv_manager, rocjpeg_calls); -} -} // namespace output -} // namespace rocpd diff --git a/projects/rocprofiler-sdk/source/lib/python/rocpd/source/csv.hpp b/projects/rocprofiler-sdk/source/lib/python/rocpd/source/csv.hpp deleted file mode 100644 index 7cabdcaf10..0000000000 --- a/projects/rocprofiler-sdk/source/lib/python/rocpd/source/csv.hpp +++ /dev/null @@ -1,120 +0,0 @@ -// MIT License -// -// Copyright (c) 2025 Advanced Micro Devices, Inc. All Rights Reserved. -// -// Permission is hereby granted, free of charge, to any person obtaining a copy -// of this software and associated documentation files (the "Software"), to deal -// in the Software without restriction, including without limitation the rights -// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -// copies of the Software, and to permit persons to whom the Software is -// furnished to do so, subject to the following conditions: -// -// The above copyright notice and this permission notice shall be included in -// all copies or substantial portions of the Software. -// -// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -// SOFTWARE. - -#pragma once - -#include "lib/python/rocpd/source/types.hpp" - -#include "lib/common/defines.hpp" -#include "lib/output/generateStats.hpp" -#include "lib/output/generator.hpp" -#include "lib/output/metadata.hpp" -#include "lib/output/node_info.hpp" -#include "lib/output/output_config.hpp" -#include "lib/output/sql/common.hpp" -#include "lib/output/stream_info.hpp" -#include "lib/rocprofiler-sdk-tool/config.hpp" - -#include -#include -#include -#include -#include -#include - -namespace rocpd -{ -namespace output -{ -using rocprofiler::tool::float_type; - -struct CsvFileConfig -{ - std::string filename; - std::string header; -}; - -enum class CsvType -{ - KERNEL_DISPATCH, - MEMORY_COPY, - MEMORY_ALLOCATION, - SCRATCH_MEMORY, - HIP_API, - HSA_CSV_API, - MARKER, - COUNTER, - RCCL_API, - ROCDECODE_API, - ROCJPEG_API, -}; - -class CsvManager -{ -public: - CsvManager(rocprofiler::tool::output_config output_cfg); - ~CsvManager(); - - rocprofiler::tool::output_config config; - std::map csv_configs; - - std::ofstream& get_stream(CsvType type); - - bool has_stream(CsvType type) const; - bool initialize_csv_file(CsvType type); - - template - void write_line(CsvType type, Args&&... args) - { - auto& stream = get_stream(type); - if(!stream.is_open()) return; - - std::vector items; - (items.push_back(fmt::format("{}", std::forward(args))), ...); - stream << fmt::format("{}\n", fmt::join(items, ",")); - } - -private: - std::map streams; - std::map file_paths; - - bool ensure_output_directory() const; -}; - -void -write_agent_info_csv(CsvManager& csv_manager, const std::vector& agents); - -void -write_csvs(CsvManager& csv_manager, - const rocprofiler::tool::generator& kernel_dispatch, - const rocprofiler::tool::generator& memory_copies, - const rocprofiler::tool::generator& memory_allocations, - const rocprofiler::tool::generator& hip_api_calls, - const rocprofiler::tool::generator& hsa_api_calls, - const rocprofiler::tool::generator& marker_api_calls, - const rocprofiler::tool::generator& counters_calls, - const rocprofiler::tool::generator& scratch_memory_calls, - const rocprofiler::tool::generator& rccl_calls, - const rocprofiler::tool::generator& rocdecode_calls, - const rocprofiler::tool::generator& rocjpeg_calls); -} // namespace output -} // namespace rocpd diff --git a/projects/rocprofiler-sdk/tests/pytest-packages/tests/rocprofv3.py b/projects/rocprofiler-sdk/tests/pytest-packages/tests/rocprofv3.py index 48f018eb0d..9cc1acfef3 100644 --- a/projects/rocprofiler-sdk/tests/pytest-packages/tests/rocprofv3.py +++ b/projects/rocprofiler-sdk/tests/pytest-packages/tests/rocprofv3.py @@ -231,3 +231,192 @@ def test_rocpd_data( assert len(_rpd_data) == len( _js_data ), f"query: {_rpd_query}\n{rpd_category} ({len(_rpd_data)}):\n\t{_rpd_data}\n{js_category} ({len(_js_data)}):\n\t{_js_data}" + + +def _perform_time_sanity_checks(data): + """Helper function to perform time sanity checks on data.""" + columns = data[0].keys() + start_columns = [c for c in columns if "start" in c.lower()] + end_columns = [c for c in columns if "end" in c.lower()] + + if not start_columns or not end_columns: + return None, None + + for record in data: + start_time = record[start_columns[0]] + end_time = record[end_columns[0]] + assert int(start_time) >= 0, f"Time error: Start time ({start_time}) < 0)." + assert int(end_time) >= 0, f"Time error: End time ({end_time}) < 0)." + assert int(end_time) >= int( + start_time + ), f"Time error: End time ({end_time}) < Start time ({start_time})." + + return start_columns[0], end_columns[0] + + +def _perform_csv_json_match(csv_row, json_row, mapping, json_data): + + def get_nested(d, path): + """Helper to get nested dict values using dot notation.""" + keys = path.split(".") + for k in keys: + if isinstance(d, dict): + d = d.get(k) + else: + return None + return d + + for csv_key, json_info in mapping.items(): + if json_info is None: + continue + + csv_value = csv_row[csv_key] + + if csv_key == "Operation": + json_value = json_data["rocprofiler-sdk-tool"]["strings"]["buffer_records"][ + json_row["kind"] + ]["operations"][json_row["operation"]] + + assert str(csv_value) in str( + json_value + ), f"Mismatch for {csv_key}: CSV={csv_value} JSON={json_value}" + continue + + if csv_key == "Function": + json_value = json_data["rocprofiler-sdk-tool"]["strings"]["buffer_records"][ + json_row["kind"] + ]["operations"][json_row["operation"]] + else: + json_path, subkey = json_info + json_value = get_nested(json_row, json_path) + if subkey: + json_value = ( + json_value.get(subkey) if isinstance(json_value, dict) else None + ) + + assert str(csv_value) == str( + json_value + ), f"Mismatch for {csv_key}: CSV={csv_value} JSON={json_value}" + + +def test_csv_data( + csv_data, + json_data, + categories=( + "agent", + "counter_collection", + "kernel", + "memory_allocation", + "memory_copy", + "regions", + ), +): + + mapping = { + "counter_collection": "counter_collection", + "kernel": "kernel_dispatch", + "memory_allocation": "memory_allocation", + "memory_copy": "memory_copy", + } + + keys_mapping = { + "kernel": { + "Thread_Id": ("thread_id", None), + "Correlation_Id": ("correlation_id", "internal"), + "Start_Timestamp": ("start_timestamp", None), + "End_Timestamp": ("end_timestamp", None), + "Queue_Id": ("dispatch_info.queue_id.handle", None), + "Kernel_Id": ("dispatch_info.kernel_id", None), + "Dispatch_Id": ("dispatch_info.dispatch_id", None), + "Stream_Id": ("stream_id.handle", None), + "Workgroup_Size_X": ("dispatch_info.workgroup_size.x", None), + "Workgroup_Size_Y": ("dispatch_info.workgroup_size.y", None), + "Workgroup_Size_Z": ("dispatch_info.workgroup_size.z", None), + "Grid_Size_X": ("dispatch_info.grid_size.x", None), + "Grid_Size_Y": ("dispatch_info.grid_size.y", None), + "Grid_Size_Z": ("dispatch_info.grid_size.z", None), + }, + "memory_allocation": { + "Operation": (), # Special case + "Correlation_Id": ("correlation_id", "internal"), + "Start_Timestamp": ("start_timestamp", None), + "End_Timestamp": ("end_timestamp", None), + }, + "memory_copy": { + "Correlation_Id": ("correlation_id", "internal"), + "Start_Timestamp": ("start_timestamp", None), + "End_Timestamp": ("end_timestamp", None), + }, + "regions": { + "Thread_Id": ("thread_id", None), + "Correlation_Id": ("correlation_id", "internal"), + "Start_Timestamp": ("start_timestamp", None), + "End_Timestamp": ("end_timestamp", None), + }, + } + + for data in csv_data: + filename, _csv_data = data + file_category = [category for category in categories if category in filename] + assert len(file_category) > 0, f"{filename} is not a valid csv filename" + category = file_category[0] + if category == "counter_collection": + _js_data = json_data["rocprofiler-sdk-tool"]["callback_records"][category] + elif category == "agent": + _js_data = json_data["rocprofiler-sdk-tool"]["agents"] + elif category == "regions": + buffer_records = json_data["rocprofiler-sdk-tool"]["buffer_records"] + _js_data = [] + for key, value in buffer_records.items(): + if key == "marker_api": + string_records = [] + marker_records = json_data["rocprofiler-sdk-tool"]["buffer_records"][ + "marker_api" + ] + for item in json_data["rocprofiler-sdk-tool"]["strings"][ + "buffer_records" + ]: + if item["kind"] == "MARKER_CORE_RANGE_API": + string_records = item["operations"] + exclude_ops = {"roctxGetThreadId"} + for entry in marker_records: + # exclude records where start and end times are the same + if entry["start_timestamp"] == entry["end_timestamp"]: + continue + # excludes roctxMarkA and roctxGetThreadId operations + if ( + string_records + and string_records[entry["operation"]] not in exclude_ops + ): + _js_data.append(entry) + else: + if key.endswith("_api"): + _js_data.extend(value) + else: + json_records_key = mapping[category] + _js_data = json_data["rocprofiler-sdk-tool"]["buffer_records"][ + json_records_key + ] + + assert len(_js_data) == len( + _csv_data + ), f"Size mismatch for {category}: JSON size= {len(_js_data)} rows, CSV size= {len(_csv_data)} rows." + + if not _csv_data: + continue # Exit if there is no data to validate + + csv_start_col, csv_end_col = _perform_time_sanity_checks(_csv_data) + json_start_col, json_end_col = _perform_time_sanity_checks(_js_data) + + if None in (csv_start_col, json_start_col, csv_end_col, json_end_col): + continue + + _csv_data_sorted = sorted( + _csv_data, key=lambda x: (int(x[csv_start_col]), int(x[csv_end_col])) + ) + _js_data_sorted = sorted( + _js_data, key=lambda x: (int(x[json_start_col]), int(x[json_end_col])) + ) + + for a, b in zip(_csv_data_sorted, _js_data_sorted): + _perform_csv_json_match(a, b, keys_mapping[category], json_data) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/CMakeLists.txt index 8d06c7599e..b413b42d7e 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/CMakeLists.txt @@ -290,7 +290,12 @@ add_test( ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --json-input ${CMAKE_CURRENT_BINARY_DIR}/rocpd-input-data/out_results.json --otf2-input ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-data/out_results.otf2 --pftrace-input - ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-data/out_results.pftrace) + ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-data/out_results.pftrace --csv-input + ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-data/out_agent_info.csv + ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-data/out_counter_collection_trace.csv + ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-data/out_kernel_trace.csv + ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-data/out_memory_allocation_trace.csv + ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-data/out_regions_trace.csv) set_tests_properties( rocprofv3-test-rocpd-validation diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/conftest.py index 73faaed82e..c35d196792 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/conftest.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/conftest.py @@ -22,6 +22,7 @@ # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN # THE SOFTWARE. +import csv import pandas as pd import pytest import json @@ -50,6 +51,12 @@ def pytest_addoption(parser): action="store", help="Path to OTF2 trace file.", ) + parser.addoption( + "--csv-input", + action="store", + nargs="+", + help="Paths to CSV files.", + ) parser.addoption( "--summary-input", action="store", @@ -83,6 +90,14 @@ def otf2_data(request): return OTF2Reader(filename).read()[0] +@pytest.fixture +def csv_data(request): + filenames = request.config.getoption("--csv-input") + return [ + (filename, list(csv.DictReader(open(filename, "r")))) for filename in filenames + ] + + @pytest.fixture def summary_data(request): filename = request.config.getoption("--summary-input") diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/validate.py index a47c56d538..86f4f979e9 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/validate.py @@ -46,6 +46,23 @@ def test_otf2_data(otf2_data, json_data): ) +def test_csv_data(csv_data, json_data): + import rocprofiler_sdk.tests.rocprofv3 as rocprofv3 + + rocprofv3.test_csv_data( + csv_data, + json_data, + ( + "agent", + "counter_collection", + "kernel", + "memory_allocation", + "memory_copy", + "regions", + ), + ) + + if __name__ == "__main__": exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) sys.exit(exit_code)