diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 0c6a5d3d9d..f8ae2f9e43 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -23,14 +23,490 @@ # THE SOFTWARE. ############################################################################### +import os + 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, query, output_path, output_file, filename="" +) -> None: + """Write the contents of a SQL query to a CSV file in the specified output path.""" + + query_one = "{} LIMIT 1".format(query) + + # just return if view is empty + if not connection.execute(query_one).fetchone(): + return + + # call query module to export to csv + file_prefix = output_file + "_" if output_file else "" + export_path = os.path.join(output_path, f"{file_prefix}{filename}_trace.csv") + export_sqlite_query(connection, query, export_format="csv", export_path=export_path) + + +def write_agent_info_csv(importData, config) -> None: + + query = """ + SELECT + guid AS Guid, + json_extract(extdata, '$.node_id') AS Node_Id, + json_extract(extdata, '$.logical_node_id') AS Logical_Node_Id, + type AS Agent_Type, + json_extract(extdata, '$.cpu_cores_count') AS Cpu_Cores_Count, + json_extract(extdata, '$.simd_count') AS Simd_Count, + json_extract(extdata, '$.cpu_core_id_base') AS Cpu_Core_Id_Base, + json_extract(extdata, '$.simd_id_base') AS Simd_Id_Base, + json_extract(extdata, '$.max_waves_per_simd') AS Max_Waves_Per_Simd, + json_extract(extdata, '$.lds_size_in_kb') AS Lds_Size_In_Kb, + json_extract(extdata, '$.gds_size_in_kb') AS Gds_Size_In_Kb, + json_extract(extdata, '$.num_gws') AS Num_Gws, + json_extract(extdata, '$.wave_front_size') AS Wave_Front_Size, + json_extract(extdata, '$.num_xcc') AS Num_Xcc, + json_extract(extdata, '$.cu_count') AS Cu_Count, + json_extract(extdata, '$.array_count') AS Array_Count, + json_extract(extdata, '$.num_shader_banks') AS Num_Shader_Banks, + json_extract(extdata, '$.simd_arrays_per_engine') AS Simd_Arrays_Per_Engine, + json_extract(extdata, '$.cu_per_simd_array') AS Cu_Per_Simd_Array, + json_extract(extdata, '$.simd_per_cu') AS Simd_Per_Cu, + json_extract(extdata, '$.max_slots_scratch_cu') AS Max_Slots_Scratch_Cu, + json_extract(extdata, '$.gfx_target_version') AS Gfx_Target_Version, + json_extract(extdata, '$.vendor_id') AS Vendor_Id, + json_extract(extdata, '$.device_id') AS Device_Id, + json_extract(extdata, '$.location_id') AS Location_Id, + json_extract(extdata, '$.domain') AS Domain, + json_extract(extdata, '$.drm_render_minor') AS Drm_Render_Minor, + json_extract(extdata, '$.num_sdma_engines') AS Num_Sdma_Engines, + json_extract(extdata, '$.num_sdma_xgmi_engines') AS Num_Sdma_Xgmi_Engines, + json_extract(extdata, '$.num_sdma_queues_per_engine') AS Num_Sdma_Queues_Per_Engine, + json_extract(extdata, '$.num_cp_queues') AS Num_Cp_Queues, + json_extract(extdata, '$.max_engine_clk_ccompute') AS Max_Engine_Clk_Ccompute, + json_extract(extdata, '$.max_engine_clk_fcompute') AS Max_Engine_Clk_Fcompute, + json_extract(extdata, '$.sdma_fw_version.uCodeSDMA') AS Sdma_Fw_Version, + json_extract(extdata, '$.fw_version.uCode') AS Fw_Version, + (COALESCE(json_extract(extdata, '$.capability.HotPluggable'), 0) << 0x0) | + (COALESCE(json_extract(extdata, '$.capability.HSAMMUPresent'), 0) << 0x1) | + (COALESCE(json_extract(extdata, '$.capability.SharedWithGraphics'), 0) << 0x2) | + (COALESCE(json_extract(extdata, '$.capability.QueueSizePowerOfTwo'), 0) << 0x3) | + (COALESCE(json_extract(extdata, '$.capability.QueueSize32bit'), 0) << 0x4) | + (COALESCE(json_extract(extdata, '$.capability.QueueIdleEvent'), 0) << 0x5) | + (COALESCE(json_extract(extdata, '$.capability.VALimit'), 0) << 0x6) | + (COALESCE(json_extract(extdata, '$.capability.WatchPointsSupported'), 0) << 0x7) | + ((COALESCE(json_extract(extdata, '$.capability.WatchPointsTotalBits'), 0) & 0xF) << 0x8) | + ((COALESCE(json_extract(extdata, '$.capability.DoorbellType'), 0) & 0x3) << 0xC) | + (COALESCE(json_extract(extdata, '$.capability.AQLQueueDoubleMap'), 0) << 0xE) | + (COALESCE(json_extract(extdata, '$.capability.DebugTrapSupported'), 0) << 0xF) | + (COALESCE(json_extract(extdata, '$.capability.WaveLaunchTrapOverrideSupported'), 0) << 0x10) | + (COALESCE(json_extract(extdata, '$.capability.WaveLaunchModeSupported'), 0) << 0x11) | + (COALESCE(json_extract(extdata, '$.capability.PreciseMemoryOperationsSupported'), 0) << 0x12) | + (COALESCE(json_extract(extdata, '$.capability.DEPRECATED_SRAM_EDCSupport'), 0) << 0x13) | + (COALESCE(json_extract(extdata, '$.capability.Mem_EDCSupport'), 0) << 0x14) | + (COALESCE(json_extract(extdata, '$.capability.RASEventNotify'), 0) << 0x15) | + ((COALESCE(json_extract(extdata, '$.capability.ASICRevision'), 0) & 0xF) << 0x16) | + (COALESCE(json_extract(extdata, '$.capability.SRAM_EDCSupport'), 0) << 0x1A) | + (COALESCE(json_extract(extdata, '$.capability.SVMAPISupported'), 0) << 0x1B) | + (COALESCE(json_extract(extdata, '$.capability.CoherentHostAccess'), 0) << 0x1C) | + (COALESCE(json_extract(extdata, '$.capability.DebugSupportedFirmware'), 0) << 0x1D) | + (COALESCE(json_extract(extdata, '$.capability.PreciseALUOperationsSupported'), 0) << 0x1E) | + (COALESCE(json_extract(extdata, '$.capability.PerQueueResetSupported'), 0) << 0x1F) AS Capability, + json_extract(extdata, '$.cu_per_engine') AS Cu_Per_Engine, + json_extract(extdata, '$.max_waves_per_cu') AS Max_Waves_Per_Cu, + json_extract(extdata, '$.workgroup_max_size') AS Workgroup_Max_Size, + json_extract(extdata, '$.family_id') AS Family_Id, + json_extract(extdata, '$.grid_max_size') AS Grid_Max_Size, + json_extract(extdata, '$.local_mem_size') AS Local_Mem_Size, + json_extract(extdata, '$.hive_id') AS Hive_Id, + json_extract(extdata, '$.gpu_id') AS Gpu_Id, + json_extract(extdata, '$.workgroup_max_dim.x') AS Workgroup_Max_Dim_X, + json_extract(extdata, '$.workgroup_max_dim.y') AS Workgroup_Max_Dim_Y, + json_extract(extdata, '$.workgroup_max_dim.z') AS Workgroup_Max_Dim_Z, + json_extract(extdata, '$.grid_max_dim.x') AS Grid_Max_Dim_X, + json_extract(extdata, '$.grid_max_dim.y') AS Grid_Max_Dim_Y, + json_extract(extdata, '$.grid_max_dim.z') AS Grid_Max_Dim_Z, + name AS Name, + json_extract(extdata, '$.vendor_name') AS Vendor_Name, + json_extract(extdata, '$.product_name') AS Product_Name, + model_name AS Model_Name + FROM "rocpd_info_agent" + """ + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "agent_info" + ) + + +def write_kernel_csv(importData, config) -> None: + + if config.agent_index_value == libpyrocpd.agent_indexing.node: # absolute + agent_id = "'Agent ' || agent_abs_index" + elif ( + config.agent_index_value == libpyrocpd.agent_indexing.logical_node + ): # relative (default) + agent_id = "'Agent ' || agent_log_index" + elif ( + config.agent_index_value == libpyrocpd.agent_indexing.logical_node_type + ): # type-relative + agent_id = "agent_type || ' ' || agent_type_index" + else: + agent_id = "" + + if config.kernel_rename: + kernel_name = "region" + else: + kernel_name = "name" + + query = f""" + SELECT + guid AS Guid, + 'KERNEL_DISPATCH' AS Kind, + {agent_id} AS Agent_Id, + queue_id AS Queue_Id, + stream_id AS Stream_Id, + tid AS Thread_Id, + dispatch_id AS Dispatch_Id, + kernel_Id AS Kernel_Id, + {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 AS Scratch_Size, + arch_vgpr_count AS Arch_VGPR_Count, + accum_vgpr_count AS Accum_VGPR_Count, + sgpr_count AS 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 + FROM "kernels" + ORDER BY + start ASC, end DESC + """ + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "kernel" + ) + + +def write_memory_copy_csv(importData, config) -> None: + + if config.agent_index_value == libpyrocpd.agent_indexing.node: # absolute + src_agent_id = "'Agent ' || src_agent_abs_index" + dst_agent_id = "'Agent ' || dst_agent_abs_index" + elif ( + config.agent_index_value == libpyrocpd.agent_indexing.logical_node + ): # relative (default) + src_agent_id = "'Agent ' || src_agent_log_index" + dst_agent_id = "'Agent ' || dst_agent_log_index" + elif ( + config.agent_index_value == libpyrocpd.agent_indexing.logical_node_type + ): # type-relative + src_agent_id = "src_agent_type || ' ' || src_agent_type_index" + dst_agent_id = "dst_agent_type || ' ' || dst_agent_type_index" + else: + src_agent_id = "" + dst_agent_id = "" + + query = f""" + SELECT + guid AS Guid, + 'MEMORY_COPY' AS Kind, + name AS Direction, + stream_id AS 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 + start ASC, end DESC + """ + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "memory_copy" + ) + + +def write_memory_allocation_csv(importData, config) -> None: + + if config.agent_index_value == libpyrocpd.agent_indexing.node: # absolute + agent_id = "'Agent ' || agent_abs_index" + elif ( + config.agent_index_value == libpyrocpd.agent_indexing.logical_node + ): # relative (default) + agent_id = "'Agent ' || agent_log_index" + elif ( + config.agent_index_value == libpyrocpd.agent_indexing.logical_node_type + ): # type-relative + agent_id = "agent_type || ' ' || agent_type_index" + else: + agent_id = "" + + query = f""" + SELECT + guid AS Guid, + 'MEMORY_ALLOCATION' AS Kind, + 'MEMORY_ALLOCATION_' || type 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 + start ASC, end DESC + """ + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "memory_allocation" + ) + + +def write_hip_api_csv(importData, config) -> None: + + query = """ + SELECT + guid AS 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" + WHERE + category LIKE 'HIP_%' + ORDER BY + start ASC, end DESC + """ + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "hip_api" + ) + + +def write_hsa_api_csv(importData, config) -> None: + + query = """ + SELECT + guid AS 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" + WHERE + category LIKE 'HSA_%' + ORDER BY + start ASC, end DESC + """ + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "hsa_api" + ) + + +def write_marker_api_csv(importData, config) -> None: + + query = """ + SELECT + guid AS Guid, + category AS Domain, + CASE + WHEN json_extract(extdata, '$.message') IS NOT NULL + THEN json_extract(extdata, '$.message') + ELSE name + END 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_and_samples" + WHERE + category LIKE 'MARKER_%' + ORDER BY + start ASC, end DESC + """ + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "marker_api" + ) + + +def write_counters_csv(importData, config) -> None: + + if config.agent_index_value == libpyrocpd.agent_indexing.node: # absolute + agent_id = "'Agent ' || agent_abs_index" + elif ( + config.agent_index_value == libpyrocpd.agent_indexing.logical_node + ): # relative (default) + agent_id = "'Agent ' || agent_log_index" + elif ( + config.agent_index_value == libpyrocpd.agent_indexing.logical_node_type + ): # type-relative + agent_id = "agent_type || ' ' || agent_type_index" + else: + agent_id = "" + + query = f""" + SELECT + guid AS Guid, + stack_id AS Correlation_Id, + dispatch_id AS Dispatch_Id, + {agent_id} AS Agent_Id, + queue_id AS Queue_Id, + pid AS Process_Id, + tid AS Thread_Id, + grid_size AS Grid_Size, + kernel_id AS Kernel_Id, + kernel_name AS Kernel_Name, + workgroup_size AS Workgroup_Size, + lds_block_size AS LDS_Block_Size, + scratch_size AS Scratch_Size, + arch_vgpr_count AS Arch_VGPR_Count, + accum_vgpr_count AS Accum_VGPR_Count, + sgpr_count AS SGPR_Count, + counter_name AS Counter_Name, + value AS Counter_Value, + start AS Start_Timestamp, + end AS End_Timestamp + FROM "counters_collection" + ORDER BY + start ASC, end DESC + """ + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "counter_collection" + ) + + +def write_scratch_memory_csv(importData, config) -> None: + + if config.agent_index_value == libpyrocpd.agent_indexing.node: # absolute + agent_id = "'Agent ' || agent_abs_index" + elif ( + config.agent_index_value == libpyrocpd.agent_indexing.logical_node + ): # relative (default) + agent_id = "'Agent ' || agent_log_index" + elif ( + config.agent_index_value == libpyrocpd.agent_indexing.logical_node_type + ): # type-relative + agent_id = "agent_type || ' ' || agent_type_index" + else: + agent_id = "" + + query = f""" + SELECT + guid AS Guid, + 'SCRATCH_MEMORY' AS Kind, + 'SCRATCH_MEMORY_' || operation AS Operation, + {agent_id} AS Agent_Id, + queue_id AS Queue_Id, + tid AS Thread_Id, + alloc_flags AS Alloc_Flags, + start AS Start_Timestamp, + end AS End_Timestamp + FROM "scratch_memory" + ORDER BY + start ASC, end DESC + """ + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "scratch_memory" + ) + + +def write_rccl_api_csv(importData, config) -> None: + + query = """ + SELECT + guid AS 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" + WHERE + category LIKE 'RCCL_%' + ORDER BY + start ASC, end DESC + """ + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "rccl_api" + ) + + +def write_rocdecode_api_csv(importData, config) -> None: + + query = """ + SELECT + guid AS 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" + WHERE + category LIKE 'ROCDECODE_%' + ORDER BY + start ASC, end DESC + """ + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "rocdecode_api" + ) + + +def write_rocjpeg_api_csv(importData, config) -> None: + + query = """ + SELECT + guid AS 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" + WHERE + category LIKE 'ROCJPEG_%' + ORDER BY + start ASC, end DESC + """ + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "rocjpeg_api" + ) + + def write_csv(importData, config): - return libpyrocpd.write_csv(importData, config) + + write_agent_info_csv(importData, config) + write_kernel_csv(importData, config) + write_memory_copy_csv(importData, config) + write_memory_allocation_csv(importData, config) + write_hip_api_csv(importData, config) + write_hsa_api_csv(importData, config) + write_marker_api_csv(importData, config) + write_counters_csv(importData, config) + write_scratch_memory_csv(importData, config) + write_rccl_api_csv(importData, config) + write_rocdecode_api_csv(importData, config) + write_rocjpeg_api_csv(importData, config) def execute(input, config=None, window_args=None, **kwargs): diff --git a/source/lib/python/rocpd/libpyrocpd.cpp b/source/lib/python/rocpd/libpyrocpd.cpp index 0f62325724..02e8e60691 100644 --- a/source/lib/python/rocpd/libpyrocpd.cpp +++ b/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/source/lib/python/rocpd/source/CMakeLists.txt b/source/lib/python/rocpd/source/CMakeLists.txt index 99a6927b43..70ac637c1f 100644 --- a/source/lib/python/rocpd/source/CMakeLists.txt +++ b/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/source/lib/python/rocpd/source/csv.cpp b/source/lib/python/rocpd/source/csv.cpp deleted file mode 100644 index 5cc5d0ac96..0000000000 --- a/source/lib/python/rocpd/source/csv.cpp +++ /dev/null @@ -1,737 +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 : 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 operation = fmt::format("MEMORY_ALLOCATION_{}", malloc.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/source/lib/python/rocpd/source/csv.hpp b/source/lib/python/rocpd/source/csv.hpp deleted file mode 100644 index 7cabdcaf10..0000000000 --- a/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/source/lib/python/rocpd/source/types.hpp b/source/lib/python/rocpd/source/types.hpp index 454ff9ec0d..f4da408ea4 100644 --- a/source/lib/python/rocpd/source/types.hpp +++ b/source/lib/python/rocpd/source/types.hpp @@ -286,7 +286,7 @@ struct kernel_dispatch uint64_t stack_id = 0; uint64_t parent_stack_id = 0; uint64_t corr_id = 0; - uint64_t vgpr_count = 0; + uint64_t arch_vgpr_count = 0; uint64_t accum_vgpr_count = 0; uint64_t sgpr_count = 0; }; @@ -444,7 +444,7 @@ struct counter uint32_t workgroup_size = 0; uint32_t lds_block_size = 0; uint32_t scratch_size = 0; - uint32_t vgpr_count = 0; + uint32_t arch_vgpr_count = 0; uint32_t accum_vgpr_count = 0; uint32_t sgpr_count = 0; std::string counter_name = {}; @@ -712,7 +712,7 @@ load(ArchiveT& ar, rocpd::types::kernel_dispatch& data) load_dim3("grid", data.grid_size); LOAD_DATA_FIELD(lds_size); LOAD_DATA_FIELD(scratch_size); - LOAD_DATA_FIELD(vgpr_count); + LOAD_DATA_FIELD(arch_vgpr_count); LOAD_DATA_FIELD(accum_vgpr_count); LOAD_DATA_FIELD(sgpr_count); LOAD_DATA_FIELD(static_lds_size); @@ -889,7 +889,7 @@ load(ArchiveT& ar, rocpd::types::counter& data) LOAD_DATA_FIELD(workgroup_size); LOAD_DATA_FIELD(lds_block_size); LOAD_DATA_FIELD(scratch_size); - LOAD_DATA_FIELD(vgpr_count); + LOAD_DATA_FIELD(arch_vgpr_count); LOAD_DATA_FIELD(accum_vgpr_count); LOAD_DATA_FIELD(sgpr_count); LOAD_DATA_FIELD(counter_name); diff --git a/source/share/rocprofiler-sdk-rocpd/data_views.sql b/source/share/rocprofiler-sdk-rocpd/data_views.sql index 0cb976688c..48b02dd8ec 100644 --- a/source/share/rocprofiler-sdk-rocpd/data_views.sql +++ b/source/share/rocprofiler-sdk-rocpd/data_views.sql @@ -315,7 +315,7 @@ SELECT K.workgroup_size_z AS workgroup_z, K.group_segment_size AS lds_size, K.private_segment_size AS scratch_size, - S.arch_vgpr_count AS vgpr_count, + S.arch_vgpr_count, S.accum_vgpr_count, S.sgpr_count, S.group_segment_size AS static_lds_size, @@ -674,7 +674,7 @@ SELECT (K.workgroup_size_x * K.workgroup_size_y * K.workgroup_size_z) AS workgroup_size, K.group_segment_size AS lds_block_size, K.private_segment_size AS scratch_size, - S.arch_vgpr_count AS vgpr_count, + S.arch_vgpr_count, S.accum_vgpr_count, S.sgpr_count, PMC_I.name AS counter_name, diff --git a/tests/pytest-packages/tests/rocprofv3.py b/tests/pytest-packages/tests/rocprofv3.py index 48f018eb0d..7143cba509 100644 --- a/tests/pytest-packages/tests/rocprofv3.py +++ b/tests/pytest-packages/tests/rocprofv3.py @@ -231,3 +231,190 @@ 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", + "hip", + "hsa", + "marker", + "kernel", + "memory_copy", + "memory_allocation", + "rocdecode_api", + "rocjpeg_api", + "counter_collection", + ), +): + + mapping = { + "hip": "hip_api", + "hsa": "hsa_api", + "marker": "marker_api", + "kernel": "kernel_dispatch", + "memory_copy": "memory_copy", + "memory_allocation": "memory_allocation", + "rocdecode_api": "rocdecode_api", + "rocjpeg_api": "rocjpeg_api", + "counter_collection": "counter_collection", + } + + 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), + }, + "hip": { + "Function": (), # Special case + "Thread_Id": ("thread_id", None), + "Correlation_Id": ("correlation_id", "internal"), + "Start_Timestamp": ("start_timestamp", None), + "End_Timestamp": ("end_timestamp", None), + }, + "hsa": { + "Function": (), # Special case + "Thread_Id": ("thread_id", None), + "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), + }, + "memory_allocation": { + "Operation": (), # Special case + "Correlation_Id": ("correlation_id", "internal"), + "Start_Timestamp": ("start_timestamp", None), + "End_Timestamp": ("end_timestamp", None), + }, + "marker": { + "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"] + 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/tests/rocprofv3/rocpd/CMakeLists.txt b/tests/rocprofv3/rocpd/CMakeLists.txt index 8d06c7599e..f4e367485b 100644 --- a/tests/rocprofv3/rocpd/CMakeLists.txt +++ b/tests/rocprofv3/rocpd/CMakeLists.txt @@ -290,7 +290,13 @@ 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_trace.csv + ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-data/out_kernel_trace.csv + ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-data/out_hip_api_trace.csv + ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-data/out_marker_api_trace.csv + ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-data/out_memory_allocation_trace.csv + ${CMAKE_CURRENT_BINARY_DIR}/rocpd-output-data/out_counter_collection_trace.csv) set_tests_properties( rocprofv3-test-rocpd-validation diff --git a/tests/rocprofv3/rocpd/conftest.py b/tests/rocprofv3/rocpd/conftest.py index 73faaed82e..c35d196792 100644 --- a/tests/rocprofv3/rocpd/conftest.py +++ b/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/tests/rocprofv3/rocpd/validate.py b/tests/rocprofv3/rocpd/validate.py index a47c56d538..4555c55374 100644 --- a/tests/rocprofv3/rocpd/validate.py +++ b/tests/rocprofv3/rocpd/validate.py @@ -46,6 +46,25 @@ 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", + "hip", + "hsa", + "marker", + "kernel", + "memory_copy", + "memory_allocation", + "counter_collection", + ), + ) + + if __name__ == "__main__": exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) sys.exit(exit_code)