From c569d146af9e8c76aeef0d22e9a633670c0c2171 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Tue, 29 Jul 2025 19:02:54 +0000 Subject: [PATCH 01/23] Write agent info to CSV --- source/lib/python/rocpd/csv.py | 104 ++++++++++++++++++++++++++++++++- 1 file changed, 103 insertions(+), 1 deletion(-) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 0c6a5d3d9d..0a7408ee04 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -23,15 +23,117 @@ # 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, filename="" +) -> None: + """Write the contents of a SQL query to a CSV file in the specified output path.""" + + # call query module to export to csv + export_path = os.path.join(output_path, f"{filename}.csv") + export_sqlite_query(connection, query, export_format="csv", export_path=export_path) + +def write_agent_info_csv( + importData, output_path +) -> None: + + query = """ + SELECT + A.guid AS Guid, + json_extract(A.extdata, '$.node_id') AS Node_Id, + json_extract(A.extdata, '$.logical_node_id') AS Logical_Node_Id, + A.type AS Agent_Type, + json_extract(A.extdata, '$.cpu_cores_count') AS Cpu_Cores_Count, + json_extract(A.extdata, '$.simd_count') AS Simd_Count, + json_extract(A.extdata, '$.cpu_core_id_base') AS Cpu_Core_Id_Base, + json_extract(A.extdata, '$.simd_id_base') AS Simd_Id_Base, + json_extract(A.extdata, '$.max_waves_per_simd') AS Max_Waves_Per_Simd, + json_extract(A.extdata, '$.lds_size_in_kb') AS Lds_Size_In_Kb, + json_extract(A.extdata, '$.gds_size_in_kb') AS Gds_Size_In_Kb, + json_extract(A.extdata, '$.num_gws') AS Num_Gws, + json_extract(A.extdata, '$.wave_front_size') AS Wave_Front_Size, + json_extract(A.extdata, '$.num_xcc') AS Num_Xcc, + json_extract(A.extdata, '$.cu_count') AS Cu_Count, + json_extract(A.extdata, '$.array_count') AS Array_Count, + json_extract(A.extdata, '$.num_shader_banks') AS Num_Shader_Banks, + json_extract(A.extdata, '$.simd_arrays_per_engine') AS Simd_Arrays_Per_Engine, + json_extract(A.extdata, '$.cu_per_simd_array') AS Cu_Per_Simd_Array, + json_extract(A.extdata, '$.simd_per_cu') AS Simd_Per_Cu, + json_extract(A.extdata, '$.max_slots_scratch_cu') AS Max_Slots_Scratch_Cu, + json_extract(A.extdata, '$.gfx_target_version') AS Gfx_Target_Version, + json_extract(A.extdata, '$.vendor_id') AS Vendor_Id, + json_extract(A.extdata, '$.device_id') AS Device_Id, + json_extract(A.extdata, '$.location_id') AS Location_Id, + json_extract(A.extdata, '$.domain') AS Domain, + json_extract(A.extdata, '$.drm_render_minor') AS Drm_Render_Minor, + json_extract(A.extdata, '$.num_sdma_engines') AS Num_Sdma_Engines, + json_extract(A.extdata, '$.num_sdma_xgmi_engines') AS Num_Sdma_Xgmi_Engines, + json_extract(A.extdata, '$.num_sdma_queues_per_engine') AS Num_Sdma_Queues_Per_Engine, + json_extract(A.extdata, '$.num_cp_queues') AS Num_Cp_Queues, + json_extract(A.extdata, '$.max_engine_clk_ccompute') AS Max_Engine_Clk_Ccompute, + json_extract(A.extdata, '$.max_engine_clk_fcompute') AS Max_Engine_Clk_Fcompute, + json_extract(A.extdata, '$.sdma_fw_version.uCodeSDMA') AS Sdma_Fw_Version, + json_extract(A.extdata, '$.fw_version.uCode') AS Fw_Version, + (COALESCE(json_extract(A.extdata, '$.capability.HotPluggable'), 0) << 0x0) | + (COALESCE(json_extract(A.extdata, '$.capability.HSAMMUPresent'), 0) << 0x1) | + (COALESCE(json_extract(A.extdata, '$.capability.SharedWithGraphics'), 0) << 0x2) | + (COALESCE(json_extract(A.extdata, '$.capability.QueueSizePowerOfTwo'), 0) << 0x3) | + (COALESCE(json_extract(A.extdata, '$.capability.QueueSize32bit'), 0) << 0x4) | + (COALESCE(json_extract(A.extdata, '$.capability.QueueIdleEvent'), 0) << 0x5) | + (COALESCE(json_extract(A.extdata, '$.capability.VALimit'), 0) << 0x6) | + (COALESCE(json_extract(A.extdata, '$.capability.WatchPointsSupported'), 0) << 0x7) | + ((COALESCE(json_extract(A.extdata, '$.capability.WatchPointsTotalBits'), 0) & 0xF) << 0x8) | + ((COALESCE(json_extract(A.extdata, '$.capability.DoorbellType'), 0) & 0x3) << 0xC) | + (COALESCE(json_extract(A.extdata, '$.capability.AQLQueueDoubleMap'), 0) << 0xE) | + (COALESCE(json_extract(A.extdata, '$.capability.DebugTrapSupported'), 0) << 0xF) | + (COALESCE(json_extract(A.extdata, '$.capability.WaveLaunchTrapOverrideSupported'), 0) << 0x10) | + (COALESCE(json_extract(A.extdata, '$.capability.WaveLaunchModeSupported'), 0) << 0x11) | + (COALESCE(json_extract(A.extdata, '$.capability.PreciseMemoryOperationsSupported'), 0) << 0x12) | + (COALESCE(json_extract(A.extdata, '$.capability.DEPRECATED_SRAM_EDCSupport'), 0) << 0x13) | + (COALESCE(json_extract(A.extdata, '$.capability.Mem_EDCSupport'), 0) << 0x14) | + (COALESCE(json_extract(A.extdata, '$.capability.RASEventNotify'), 0) << 0x15) | + ((COALESCE(json_extract(A.extdata, '$.capability.ASICRevision'), 0) & 0xF) << 0x16) | + (COALESCE(json_extract(A.extdata, '$.capability.SRAM_EDCSupport'), 0) << 0x1A) | + (COALESCE(json_extract(A.extdata, '$.capability.SVMAPISupported'), 0) << 0x1B) | + (COALESCE(json_extract(A.extdata, '$.capability.CoherentHostAccess'), 0) << 0x1C) | + (COALESCE(json_extract(A.extdata, '$.capability.DebugSupportedFirmware'), 0) << 0x1D) | + (COALESCE(json_extract(A.extdata, '$.capability.PreciseALUOperationsSupported'), 0) << 0x1E) | + (COALESCE(json_extract(A.extdata, '$.capability.PerQueueResetSupported'), 0) << 0x1F) AS Capability, + json_extract(A.extdata, '$.cu_per_engine') AS Cu_Per_Engine, + json_extract(A.extdata, '$.max_waves_per_cu') AS Max_Waves_Per_Cu, + json_extract(A.extdata, '$.workgroup_max_size') AS Workgroup_Max_Size, + json_extract(A.extdata, '$.family_id') AS Family_Id, + json_extract(A.extdata, '$.grid_max_size') AS Grid_Max_Size, + json_extract(A.extdata, '$.local_mem_size') AS Local_Mem_Size, + json_extract(A.extdata, '$.hive_id') AS Hive_Id, + json_extract(A.extdata, '$.gpu_id') AS Gpu_Id, + json_extract(A.extdata, '$.workgroup_max_dim.x') AS Workgroup_Max_Dim_X, + json_extract(A.extdata, '$.workgroup_max_dim.y') AS Workgroup_Max_Dim_Y, + json_extract(A.extdata, '$.workgroup_max_dim.z') AS Workgroup_Max_Dim_Z, + json_extract(A.extdata, '$.grid_max_dim.x') AS Grid_Max_Dim_X, + json_extract(A.extdata, '$.grid_max_dim.y') AS Grid_Max_Dim_Y, + json_extract(A.extdata, '$.grid_max_dim.z') AS Grid_Max_Dim_Z, + A.name AS Name, + json_extract(A.extdata, '$.vendor_name') AS Vendor_Name, + json_extract(A.extdata, '$.product_name') AS Product_Name, + A.model_name AS Model_Name + FROM "rocpd_info_node" AS N + INNER JOIN rocpd_info_agent as A + ON A.guid = N.guid + AND A.nid = N.id + """ + write_sql_query_to_csv(importData, query, output_path, "out_agent_info") def write_csv(importData, config): - return libpyrocpd.write_csv(importData, config) + write_agent_info_csv(importData, config.output_path) def execute(input, config=None, window_args=None, **kwargs): From 813e0c9d54f6b2ad530dcbac2ce83da621841394 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Wed, 30 Jul 2025 13:04:50 +0000 Subject: [PATCH 02/23] Write kernel to CSV --- source/lib/python/rocpd/csv.py | 40 ++++++++++++++++++++++++++++++++++ 1 file changed, 40 insertions(+) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 0a7408ee04..0f0acd95b4 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -131,9 +131,49 @@ def write_agent_info_csv( """ write_sql_query_to_csv(importData, query, output_path, "out_agent_info") +def write_kernel_csv( + importData, output_path +) -> None: + + query = """ + SELECT + K.guid AS Guid, + 'KERNEL_DISPATCH' AS Kind, + 'Agent ' || K.agent_log_index AS Agent_Id, + K.queue_id AS Queue_Id, + K.stream_id AS Stream_Id, + K.tid AS Thread_Id, + K.dispatch_id AS Dispatch_Id, + K.kernel_Id AS Kernel_Id, + K.name AS Kernel_Name, + K.stack_id AS Correlation_Id, + K.start AS Start_Timestamp, + K.end AS End_Timestamp, + K.scratch_size AS Private_Segment_Size, + K.lds_size AS Group_Segment_Size, + K.workgroup_x AS Workgroup_Size_X, + K.workgroup_y AS Workgroup_Size_Y, + K.workgroup_z AS Workgroup_Size_Z, + K.grid_x AS Grid_Size_X, + K.grid_y AS Grid_Size_Y, + K.grid_z AS Grid_Size_Z + FROM "rocpd_info_node" AS N + INNER JOIN rocpd_info_process as P + ON P.guid = N.guid + AND P.nid = N.id + INNER JOIN kernels AS K + ON K.guid = P.guid + AND K.nid = P.nid + AND K.pid = P.pid + ORDER BY + K.start ASC, K.end DESC + """ + write_sql_query_to_csv(importData, query, output_path, "out_kernel_trace") + def write_csv(importData, config): write_agent_info_csv(importData, config.output_path) + write_kernel_csv(importData, config.output_path) def execute(input, config=None, window_args=None, **kwargs): From 44111b5cad2d9ea1c1b7b84eb544e6ed8b967cca Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Wed, 30 Jul 2025 14:26:09 +0000 Subject: [PATCH 03/23] Write memory copy to CSV --- source/lib/python/rocpd/csv.py | 30 ++++++++++++++++++++++++++++++ 1 file changed, 30 insertions(+) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 0f0acd95b4..ca1d033ec5 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -170,10 +170,40 @@ def write_kernel_csv( """ write_sql_query_to_csv(importData, query, output_path, "out_kernel_trace") +def write_memory_copy_csv( + importData, output_path +) -> None: + + query = """ + SELECT + M.guid AS Guid, + 'MEMORY_COPY' AS Kind, + M.name AS Direction, + M.stream_id AS Stream_Id, + 'Agent ' || M.src_agent_log_index AS Source_Agent_Id, + 'Agent ' || M.dst_agent_log_index AS Destination_Agent_Id, + M.stack_id AS Correlation_Id, + M.start AS Start_Timestamp, + M.end AS End_Timestamp + + FROM "rocpd_info_node" AS N + INNER JOIN rocpd_info_process as P + ON P.guid = N.guid + AND P.nid = N.id + INNER JOIN memory_copies AS M + ON M.guid = P.guid + AND M.nid = P.nid + AND M.pid = P.pid + ORDER BY + M.start ASC, M.end DESC + """ + write_sql_query_to_csv(importData, query, output_path, "out_memory_copy_trace") + def write_csv(importData, config): write_agent_info_csv(importData, config.output_path) write_kernel_csv(importData, config.output_path) + write_memory_copy_csv(importData, config.output_path) def execute(input, config=None, window_args=None, **kwargs): From 24c3f852a6ec796198fa0c37817ff83322535a95 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Wed, 30 Jul 2025 14:36:45 +0000 Subject: [PATCH 04/23] Write memory allocation to CSV --- source/lib/python/rocpd/csv.py | 33 +++++++++++++++++++++++++++++++++ 1 file changed, 33 insertions(+) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index ca1d033ec5..b8c9abadb4 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -199,11 +199,44 @@ def write_memory_copy_csv( """ write_sql_query_to_csv(importData, query, output_path, "out_memory_copy_trace") +def write_memory_allocation_csv( + importData, output_path +) -> None: + + query = """ + SELECT + A.guid AS Guid, + 'MEMORY_ALLOCATION' AS Kind, + 'MEMORY_ALLOCATION_' || A.type AS Operation, + CASE + WHEN A.type != "FREE" + THEN 'Agent ' || A.agent_log_index + ELSE '"' + END AS Agent_Id, + A.size AS Allocation_Size, + '0x' || printf('%016X', A.address) AS Address, + A.stack_id AS Correlation_Id, + A.start AS Start_Timestamp, + A.end AS End_Timestamp + FROM "rocpd_info_node" AS N + INNER JOIN rocpd_info_process as P + ON P.guid = N.guid + AND P.nid = N.id + INNER JOIN memory_allocations AS A + ON A.guid = P.guid + AND A.nid = P.nid + AND A.pid = P.pid + ORDER BY + A.start ASC, A.end DESC + """ + write_sql_query_to_csv(importData, query, output_path, "out_memory_allocation_trace") + def write_csv(importData, config): write_agent_info_csv(importData, config.output_path) write_kernel_csv(importData, config.output_path) write_memory_copy_csv(importData, config.output_path) + write_memory_allocation_csv(importData, config.output_path) def execute(input, config=None, window_args=None, **kwargs): From 9cb4c01c33fc70de309f5294a3af0776a86044bd Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Wed, 30 Jul 2025 14:49:54 +0000 Subject: [PATCH 05/23] Write hip api to CSV --- source/lib/python/rocpd/csv.py | 30 ++++++++++++++++++++++++++++++ 1 file changed, 30 insertions(+) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index b8c9abadb4..934a6865b8 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -231,12 +231,42 @@ def write_memory_allocation_csv( """ write_sql_query_to_csv(importData, query, output_path, "out_memory_allocation_trace") +def write_hip_api_csv( + importData, output_path +) -> None: + + query = """ + SELECT + R.guid AS Guid, + R.category AS Domain, + R.name AS Function, + R.pid AS Process_Id, + R.tid AS Thread_Id, + R.stack_id AS Correlation_Id, + R.start AS Start_Timestamp, + R.end AS End_Timestamp + FROM "rocpd_info_node" AS N + INNER JOIN rocpd_info_process as P + ON P.guid = N.guid + AND P.nid = N.id + INNER JOIN regions AS R + ON R.guid = P.guid + AND R.nid = P.nid + AND R.pid = P.pid + WHERE + R.category LIKE 'HIP_%' + ORDER BY + R.start ASC, R.end DESC + """ + write_sql_query_to_csv(importData, query, output_path, "out_hip_api_trace") + def write_csv(importData, config): write_agent_info_csv(importData, config.output_path) write_kernel_csv(importData, config.output_path) write_memory_copy_csv(importData, config.output_path) write_memory_allocation_csv(importData, config.output_path) + write_hip_api_csv(importData, config.output_path) def execute(input, config=None, window_args=None, **kwargs): From b031cfe0a3d92d61118124b2ed389e91ece8bc9a Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Wed, 30 Jul 2025 14:55:49 +0000 Subject: [PATCH 06/23] Write hsa api to CSV --- source/lib/python/rocpd/csv.py | 30 ++++++++++++++++++++++++++++++ 1 file changed, 30 insertions(+) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 934a6865b8..be59db3161 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -260,6 +260,35 @@ def write_hip_api_csv( """ write_sql_query_to_csv(importData, query, output_path, "out_hip_api_trace") +def write_hsa_api_csv( + importData, output_path +) -> None: + + query = """ + SELECT + R.guid AS Guid, + R.category AS Domain, + R.name AS Function, + R.pid AS Process_Id, + R.tid AS Thread_Id, + R.stack_id AS Correlation_Id, + R.start AS Start_Timestamp, + R.end AS End_Timestamp + FROM "rocpd_info_node" AS N + INNER JOIN rocpd_info_process as P + ON P.guid = N.guid + AND P.nid = N.id + INNER JOIN regions AS R + ON R.guid = P.guid + AND R.nid = P.nid + AND R.pid = P.pid + WHERE + R.category LIKE 'HSA_%' + ORDER BY + R.start ASC, R.end DESC + """ + write_sql_query_to_csv(importData, query, output_path, "out_hsa_api_trace") + def write_csv(importData, config): write_agent_info_csv(importData, config.output_path) @@ -267,6 +296,7 @@ def write_csv(importData, config): write_memory_copy_csv(importData, config.output_path) write_memory_allocation_csv(importData, config.output_path) write_hip_api_csv(importData, config.output_path) + write_hsa_api_csv(importData, config.output_path) def execute(input, config=None, window_args=None, **kwargs): From 81a2996ce0accfdfd79d236b7d9c009c5a248150 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Wed, 30 Jul 2025 17:36:54 +0000 Subject: [PATCH 07/23] Write marker api to CSV --- source/lib/python/rocpd/csv.py | 34 ++++++++++++++++++++++++++++++++++ 1 file changed, 34 insertions(+) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index be59db3161..1ac14a2cce 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -289,6 +289,39 @@ def write_hsa_api_csv( """ write_sql_query_to_csv(importData, query, output_path, "out_hsa_api_trace") +def write_marker_api_csv( + importData, output_path +) -> None: + + query = """ + SELECT + R.guid AS Guid, + R.category AS Domain, + CASE + WHEN json_extract(R.extdata, '$.message') IS NOT NULL + THEN json_extract(R.extdata, '$.message') + ELSE R.name + END AS Function, + R.pid AS Process_Id, + R.tid AS Thread_Id, + R.stack_id AS Correlation_Id, + R.start AS Start_Timestamp, + R.end AS End_Timestamp + FROM "rocpd_info_node" AS N + INNER JOIN rocpd_info_process as P + ON P.guid = N.guid + AND P.nid = N.id + INNER JOIN regions_and_samples AS R + ON R.guid = P.guid + AND R.nid = P.nid + AND R.pid = P.pid + WHERE + R.category LIKE 'MARKER_%' + ORDER BY + R.start ASC, R.end DESC + """ + write_sql_query_to_csv(importData, query, output_path, "out_marker_api_trace") + def write_csv(importData, config): write_agent_info_csv(importData, config.output_path) @@ -297,6 +330,7 @@ def write_csv(importData, config): write_memory_allocation_csv(importData, config.output_path) write_hip_api_csv(importData, config.output_path) write_hsa_api_csv(importData, config.output_path) + write_marker_api_csv(importData, config.output_path) def execute(input, config=None, window_args=None, **kwargs): From 9ec960cf04c6247ec95c22f35faa38d1894a72cd Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Thu, 31 Jul 2025 15:57:20 +0000 Subject: [PATCH 08/23] Write counters to CSV --- source/lib/python/rocpd/csv.py | 40 ++++++++++++++++++++++++++++++++++ 1 file changed, 40 insertions(+) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 1ac14a2cce..1ac2b8b604 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -322,6 +322,45 @@ def write_marker_api_csv( """ write_sql_query_to_csv(importData, query, output_path, "out_marker_api_trace") +def write_counters_csv( + importData, output_path +) -> None: + + query = """ + SELECT + C.guid AS Pid, + C.stack_id AS Correlation_Id, + C.dispatch_id AS Dispatch_Id, + 'Agent ' || C.agent_log_index AS Agent_Id, + C.queue_id AS Queue_Id, + C.pid AS Process_Id, + C.tid AS Thread_Id, + C.grid_size AS Grid_Size, + C.kernel_id AS Kernel_Id, + C.kernel_name AS Kernel_Name, + C.workgroup_size AS Workgroup_Size, + C.lds_block_size AS LDS_Block_Size, + C.scratch_size AS Scratch_Size, + C.vgpr_count AS VGPR_Count, + C.accum_vgpr_count AS Accum_VGPR_Count, + C.sgpr_count AS SGPR_Count, + C.counter_name AS Counter_Name, + C.value AS Counter_Value, + C.start AS Start_Timestamp, + C.end AS End_Timestamp + FROM "rocpd_info_node" AS N + INNER JOIN rocpd_info_process as P + ON P.guid = N.guid + AND P.nid = N.id + INNER JOIN counters_collection AS C + ON C.guid = P.guid + AND C.nid = P.nid + AND C.pid = P.pid + ORDER BY + C.start ASC, C.end DESC + """ + write_sql_query_to_csv(importData, query, output_path, "out_counter_collection") + def write_csv(importData, config): write_agent_info_csv(importData, config.output_path) @@ -331,6 +370,7 @@ def write_csv(importData, config): write_hip_api_csv(importData, config.output_path) write_hsa_api_csv(importData, config.output_path) write_marker_api_csv(importData, config.output_path) + write_counters_csv(importData, config.output_path) def execute(input, config=None, window_args=None, **kwargs): From 255e3b144a52f636a719af086dce16436ae227f6 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Thu, 31 Jul 2025 15:58:40 +0000 Subject: [PATCH 09/23] Write scratch memory to CSV --- source/lib/python/rocpd/csv.py | 28 ++++++++++++++++++++++++++++ 1 file changed, 28 insertions(+) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 1ac2b8b604..551b2e258a 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -361,6 +361,33 @@ def write_counters_csv( """ write_sql_query_to_csv(importData, query, output_path, "out_counter_collection") +def write_scratch_memory_csv( + importData, output_path +) -> None: + + query = """ + SELECT + 'SCRATCH_MEMORY' AS Kind, + 'SCRATCH_MEMORY_' || S.operation AS Operation, + 'Agent ' || S.agent_log_index AS Agent_Id, + S.queue_id AS Queue_Id, + S.tid AS Thread_Id, + S.alloc_flags AS Alloc_Flags, + S.start AS Start_Timestamp, + S.end AS End_Timestamp + FROM "rocpd_info_node" AS N + INNER JOIN rocpd_info_process as P + ON P.guid = N.guid + AND P.nid = N.id + INNER JOIN scratch_memory AS S + ON S.guid = P.guid + AND S.nid = P.nid + AND S.pid = P.pid + ORDER BY + S.start ASC, S.end + """ + write_sql_query_to_csv(importData, query, output_path, "out_scratch_memory_trace") + def write_csv(importData, config): write_agent_info_csv(importData, config.output_path) @@ -371,6 +398,7 @@ def write_csv(importData, config): write_hsa_api_csv(importData, config.output_path) write_marker_api_csv(importData, config.output_path) write_counters_csv(importData, config.output_path) + write_scratch_memory_csv(importData, config.output_path) def execute(input, config=None, window_args=None, **kwargs): From dde982ce131530bfc8089c991a3d23b5ac2de106 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Thu, 31 Jul 2025 15:59:59 +0000 Subject: [PATCH 10/23] Write rccl api to CSV --- source/lib/python/rocpd/csv.py | 31 +++++++++++++++++++++++++++++++ 1 file changed, 31 insertions(+) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 551b2e258a..c9a46773a0 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -388,6 +388,36 @@ def write_scratch_memory_csv( """ write_sql_query_to_csv(importData, query, output_path, "out_scratch_memory_trace") +def write_rccl_api_csv( + importData, output_path +) -> None: + + query = """ + SELECT + R.guid AS Guid, + R.category AS Domain, + R.name AS Function, + R.pid AS Process_Id, + R.tid AS Thread_Id, + R.stack_id AS Correlation_Id, + R.start AS Start_Timestamp, + R.end AS End_Timestamp + FROM "rocpd_info_node" AS N + INNER JOIN rocpd_info_process as P + ON P.guid = N.guid + AND P.nid = N.id + INNER JOIN regions AS R + ON R.guid = P.guid + AND R.nid = P.nid + AND R.pid = P.pid + WHERE + R.category LIKE 'RCCL_%' + ORDER BY + R.start ASC, R.end + """ + write_sql_query_to_csv(importData, query, output_path, "out_rccl_api_trace") + + def write_csv(importData, config): write_agent_info_csv(importData, config.output_path) @@ -399,6 +429,7 @@ def write_csv(importData, config): write_marker_api_csv(importData, config.output_path) write_counters_csv(importData, config.output_path) write_scratch_memory_csv(importData, config.output_path) + write_rccl_api_csv(importData, config.output_path) def execute(input, config=None, window_args=None, **kwargs): From e69fc43cf732542c252b724782b5e3d0ecc711aa Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Thu, 31 Jul 2025 16:01:17 +0000 Subject: [PATCH 11/23] Write rocdecode api to CSV --- source/lib/python/rocpd/csv.py | 29 +++++++++++++++++++++++++++++ 1 file changed, 29 insertions(+) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index c9a46773a0..7ce0c7f226 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -417,6 +417,34 @@ def write_rccl_api_csv( """ write_sql_query_to_csv(importData, query, output_path, "out_rccl_api_trace") +def write_rocdecode_api_csv( + importData, output_path +) -> None: + + query = """ + SELECT + R.guid AS Guid, + R.category AS Domain, + R.name AS Function, + R.pid AS Process_Id, + R.tid AS Thread_Id, + R.stack_id AS Correlation_Id, + R.start AS Start_Timestamp, + R.end AS End_Timestamp + FROM "rocpd_info_node" AS N + INNER JOIN rocpd_info_process as P + ON P.guid = N.guid + AND P.nid = N.id + INNER JOIN regions AS R + ON R.guid = P.guid + AND R.nid = P.nid + AND R.pid = P.pid + WHERE + R.category LIKE 'ROCDECODE_%' + ORDER BY + R.start ASC, R.end + """ + write_sql_query_to_csv(importData, query, output_path, "out_rocdecode_api_trace") def write_csv(importData, config): @@ -430,6 +458,7 @@ def write_csv(importData, config): write_counters_csv(importData, config.output_path) write_scratch_memory_csv(importData, config.output_path) write_rccl_api_csv(importData, config.output_path) + write_rocdecode_api_csv(importData, config.output_path) def execute(input, config=None, window_args=None, **kwargs): From efc6564ea7878dbd37d63a8cdd22612c9040dab8 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Thu, 31 Jul 2025 16:04:41 +0000 Subject: [PATCH 12/23] Write rocjpeg api to CSV --- source/lib/python/rocpd/csv.py | 30 ++++++++++++++++++++++++++++++ 1 file changed, 30 insertions(+) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 7ce0c7f226..3d418d5e14 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -446,6 +446,35 @@ def write_rocdecode_api_csv( """ write_sql_query_to_csv(importData, query, output_path, "out_rocdecode_api_trace") +def write_rocjpeg_api_csv( + importData, output_path +) -> None: + + query = """ + SELECT + R.guid AS Guid, + R.category AS Domain, + R.name AS Function, + R.pid AS Process_Id, + R.tid AS Thread_Id, + R.stack_id AS Correlation_Id, + R.start AS Start_Timestamp, + R.end AS End_Timestamp + FROM "rocpd_info_node" AS N + INNER JOIN rocpd_info_process as P + ON P.guid = N.guid + AND P.nid = N.id + INNER JOIN regions AS R + ON R.guid = P.guid + AND R.nid = P.nid + AND R.pid = P.pid + WHERE + R.category LIKE 'ROCJPEG_%' + ORDER BY + R.start ASC, R.end + """ + write_sql_query_to_csv(importData, query, output_path, "out_rocjpeg_api_trace") + def write_csv(importData, config): write_agent_info_csv(importData, config.output_path) @@ -459,6 +488,7 @@ def write_csv(importData, config): write_scratch_memory_csv(importData, config.output_path) write_rccl_api_csv(importData, config.output_path) write_rocdecode_api_csv(importData, config.output_path) + write_rocjpeg_api_csv(importData, config.output_path) def execute(input, config=None, window_args=None, **kwargs): From 42bba0056a094c9103eb219b67cdefd366a6b023 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Thu, 31 Jul 2025 16:22:33 +0000 Subject: [PATCH 13/23] Remove info_process joins --- source/lib/python/rocpd/csv.py | 521 ++++++++++++++------------------- 1 file changed, 220 insertions(+), 301 deletions(-) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 3d418d5e14..7a1737fd34 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -46,88 +46,85 @@ def write_agent_info_csv( query = """ SELECT - A.guid AS Guid, - json_extract(A.extdata, '$.node_id') AS Node_Id, - json_extract(A.extdata, '$.logical_node_id') AS Logical_Node_Id, - A.type AS Agent_Type, - json_extract(A.extdata, '$.cpu_cores_count') AS Cpu_Cores_Count, - json_extract(A.extdata, '$.simd_count') AS Simd_Count, - json_extract(A.extdata, '$.cpu_core_id_base') AS Cpu_Core_Id_Base, - json_extract(A.extdata, '$.simd_id_base') AS Simd_Id_Base, - json_extract(A.extdata, '$.max_waves_per_simd') AS Max_Waves_Per_Simd, - json_extract(A.extdata, '$.lds_size_in_kb') AS Lds_Size_In_Kb, - json_extract(A.extdata, '$.gds_size_in_kb') AS Gds_Size_In_Kb, - json_extract(A.extdata, '$.num_gws') AS Num_Gws, - json_extract(A.extdata, '$.wave_front_size') AS Wave_Front_Size, - json_extract(A.extdata, '$.num_xcc') AS Num_Xcc, - json_extract(A.extdata, '$.cu_count') AS Cu_Count, - json_extract(A.extdata, '$.array_count') AS Array_Count, - json_extract(A.extdata, '$.num_shader_banks') AS Num_Shader_Banks, - json_extract(A.extdata, '$.simd_arrays_per_engine') AS Simd_Arrays_Per_Engine, - json_extract(A.extdata, '$.cu_per_simd_array') AS Cu_Per_Simd_Array, - json_extract(A.extdata, '$.simd_per_cu') AS Simd_Per_Cu, - json_extract(A.extdata, '$.max_slots_scratch_cu') AS Max_Slots_Scratch_Cu, - json_extract(A.extdata, '$.gfx_target_version') AS Gfx_Target_Version, - json_extract(A.extdata, '$.vendor_id') AS Vendor_Id, - json_extract(A.extdata, '$.device_id') AS Device_Id, - json_extract(A.extdata, '$.location_id') AS Location_Id, - json_extract(A.extdata, '$.domain') AS Domain, - json_extract(A.extdata, '$.drm_render_minor') AS Drm_Render_Minor, - json_extract(A.extdata, '$.num_sdma_engines') AS Num_Sdma_Engines, - json_extract(A.extdata, '$.num_sdma_xgmi_engines') AS Num_Sdma_Xgmi_Engines, - json_extract(A.extdata, '$.num_sdma_queues_per_engine') AS Num_Sdma_Queues_Per_Engine, - json_extract(A.extdata, '$.num_cp_queues') AS Num_Cp_Queues, - json_extract(A.extdata, '$.max_engine_clk_ccompute') AS Max_Engine_Clk_Ccompute, - json_extract(A.extdata, '$.max_engine_clk_fcompute') AS Max_Engine_Clk_Fcompute, - json_extract(A.extdata, '$.sdma_fw_version.uCodeSDMA') AS Sdma_Fw_Version, - json_extract(A.extdata, '$.fw_version.uCode') AS Fw_Version, - (COALESCE(json_extract(A.extdata, '$.capability.HotPluggable'), 0) << 0x0) | - (COALESCE(json_extract(A.extdata, '$.capability.HSAMMUPresent'), 0) << 0x1) | - (COALESCE(json_extract(A.extdata, '$.capability.SharedWithGraphics'), 0) << 0x2) | - (COALESCE(json_extract(A.extdata, '$.capability.QueueSizePowerOfTwo'), 0) << 0x3) | - (COALESCE(json_extract(A.extdata, '$.capability.QueueSize32bit'), 0) << 0x4) | - (COALESCE(json_extract(A.extdata, '$.capability.QueueIdleEvent'), 0) << 0x5) | - (COALESCE(json_extract(A.extdata, '$.capability.VALimit'), 0) << 0x6) | - (COALESCE(json_extract(A.extdata, '$.capability.WatchPointsSupported'), 0) << 0x7) | - ((COALESCE(json_extract(A.extdata, '$.capability.WatchPointsTotalBits'), 0) & 0xF) << 0x8) | - ((COALESCE(json_extract(A.extdata, '$.capability.DoorbellType'), 0) & 0x3) << 0xC) | - (COALESCE(json_extract(A.extdata, '$.capability.AQLQueueDoubleMap'), 0) << 0xE) | - (COALESCE(json_extract(A.extdata, '$.capability.DebugTrapSupported'), 0) << 0xF) | - (COALESCE(json_extract(A.extdata, '$.capability.WaveLaunchTrapOverrideSupported'), 0) << 0x10) | - (COALESCE(json_extract(A.extdata, '$.capability.WaveLaunchModeSupported'), 0) << 0x11) | - (COALESCE(json_extract(A.extdata, '$.capability.PreciseMemoryOperationsSupported'), 0) << 0x12) | - (COALESCE(json_extract(A.extdata, '$.capability.DEPRECATED_SRAM_EDCSupport'), 0) << 0x13) | - (COALESCE(json_extract(A.extdata, '$.capability.Mem_EDCSupport'), 0) << 0x14) | - (COALESCE(json_extract(A.extdata, '$.capability.RASEventNotify'), 0) << 0x15) | - ((COALESCE(json_extract(A.extdata, '$.capability.ASICRevision'), 0) & 0xF) << 0x16) | - (COALESCE(json_extract(A.extdata, '$.capability.SRAM_EDCSupport'), 0) << 0x1A) | - (COALESCE(json_extract(A.extdata, '$.capability.SVMAPISupported'), 0) << 0x1B) | - (COALESCE(json_extract(A.extdata, '$.capability.CoherentHostAccess'), 0) << 0x1C) | - (COALESCE(json_extract(A.extdata, '$.capability.DebugSupportedFirmware'), 0) << 0x1D) | - (COALESCE(json_extract(A.extdata, '$.capability.PreciseALUOperationsSupported'), 0) << 0x1E) | - (COALESCE(json_extract(A.extdata, '$.capability.PerQueueResetSupported'), 0) << 0x1F) AS Capability, - json_extract(A.extdata, '$.cu_per_engine') AS Cu_Per_Engine, - json_extract(A.extdata, '$.max_waves_per_cu') AS Max_Waves_Per_Cu, - json_extract(A.extdata, '$.workgroup_max_size') AS Workgroup_Max_Size, - json_extract(A.extdata, '$.family_id') AS Family_Id, - json_extract(A.extdata, '$.grid_max_size') AS Grid_Max_Size, - json_extract(A.extdata, '$.local_mem_size') AS Local_Mem_Size, - json_extract(A.extdata, '$.hive_id') AS Hive_Id, - json_extract(A.extdata, '$.gpu_id') AS Gpu_Id, - json_extract(A.extdata, '$.workgroup_max_dim.x') AS Workgroup_Max_Dim_X, - json_extract(A.extdata, '$.workgroup_max_dim.y') AS Workgroup_Max_Dim_Y, - json_extract(A.extdata, '$.workgroup_max_dim.z') AS Workgroup_Max_Dim_Z, - json_extract(A.extdata, '$.grid_max_dim.x') AS Grid_Max_Dim_X, - json_extract(A.extdata, '$.grid_max_dim.y') AS Grid_Max_Dim_Y, - json_extract(A.extdata, '$.grid_max_dim.z') AS Grid_Max_Dim_Z, - A.name AS Name, - json_extract(A.extdata, '$.vendor_name') AS Vendor_Name, - json_extract(A.extdata, '$.product_name') AS Product_Name, - A.model_name AS Model_Name - FROM "rocpd_info_node" AS N - INNER JOIN rocpd_info_agent as A - ON A.guid = N.guid - AND A.nid = N.id + 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, output_path, "out_agent_info") @@ -137,36 +134,29 @@ def write_kernel_csv( query = """ SELECT - K.guid AS Guid, + guid AS Guid, 'KERNEL_DISPATCH' AS Kind, - 'Agent ' || K.agent_log_index AS Agent_Id, - K.queue_id AS Queue_Id, - K.stream_id AS Stream_Id, - K.tid AS Thread_Id, - K.dispatch_id AS Dispatch_Id, - K.kernel_Id AS Kernel_Id, - K.name AS Kernel_Name, - K.stack_id AS Correlation_Id, - K.start AS Start_Timestamp, - K.end AS End_Timestamp, - K.scratch_size AS Private_Segment_Size, - K.lds_size AS Group_Segment_Size, - K.workgroup_x AS Workgroup_Size_X, - K.workgroup_y AS Workgroup_Size_Y, - K.workgroup_z AS Workgroup_Size_Z, - K.grid_x AS Grid_Size_X, - K.grid_y AS Grid_Size_Y, - K.grid_z AS Grid_Size_Z - FROM "rocpd_info_node" AS N - INNER JOIN rocpd_info_process as P - ON P.guid = N.guid - AND P.nid = N.id - INNER JOIN kernels AS K - ON K.guid = P.guid - AND K.nid = P.nid - AND K.pid = P.pid + 'Agent ' || agent_log_index 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, + name AS Kernel_Name, + stack_id AS Correlation_Id, + start AS Start_Timestamp, + end AS End_Timestamp, + scratch_size AS Private_Segment_Size, + lds_size AS Group_Segment_Size, + 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 - K.start ASC, K.end DESC + start ASC, end DESC """ write_sql_query_to_csv(importData, query, output_path, "out_kernel_trace") @@ -176,26 +166,18 @@ def write_memory_copy_csv( query = """ SELECT - M.guid AS Guid, + guid AS Guid, 'MEMORY_COPY' AS Kind, - M.name AS Direction, - M.stream_id AS Stream_Id, - 'Agent ' || M.src_agent_log_index AS Source_Agent_Id, - 'Agent ' || M.dst_agent_log_index AS Destination_Agent_Id, - M.stack_id AS Correlation_Id, - M.start AS Start_Timestamp, - M.end AS End_Timestamp - - FROM "rocpd_info_node" AS N - INNER JOIN rocpd_info_process as P - ON P.guid = N.guid - AND P.nid = N.id - INNER JOIN memory_copies AS M - ON M.guid = P.guid - AND M.nid = P.nid - AND M.pid = P.pid + name AS Direction, + stream_id AS Stream_Id, + 'Agent ' || src_agent_log_index AS Source_Agent_Id, + 'Agent ' || dst_agent_log_index AS Destination_Agent_Id, + stack_id AS Correlation_Id, + start AS Start_Timestamp, + end AS End_Timestamp + FROM "memory_copies" ORDER BY - M.start ASC, M.end DESC + start ASC, end DESC """ write_sql_query_to_csv(importData, query, output_path, "out_memory_copy_trace") @@ -205,29 +187,22 @@ def write_memory_allocation_csv( query = """ SELECT - A.guid AS Guid, + guid AS Guid, 'MEMORY_ALLOCATION' AS Kind, - 'MEMORY_ALLOCATION_' || A.type AS Operation, + 'MEMORY_ALLOCATION_' || type AS Operation, CASE - WHEN A.type != "FREE" - THEN 'Agent ' || A.agent_log_index + WHEN type != "FREE" + THEN 'Agent ' || agent_log_index ELSE '"' END AS Agent_Id, - A.size AS Allocation_Size, - '0x' || printf('%016X', A.address) AS Address, - A.stack_id AS Correlation_Id, - A.start AS Start_Timestamp, - A.end AS End_Timestamp - FROM "rocpd_info_node" AS N - INNER JOIN rocpd_info_process as P - ON P.guid = N.guid - AND P.nid = N.id - INNER JOIN memory_allocations AS A - ON A.guid = P.guid - AND A.nid = P.nid - AND A.pid = P.pid + 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 - A.start ASC, A.end DESC + start ASC, end DESC """ write_sql_query_to_csv(importData, query, output_path, "out_memory_allocation_trace") @@ -237,26 +212,19 @@ def write_hip_api_csv( query = """ SELECT - R.guid AS Guid, - R.category AS Domain, - R.name AS Function, - R.pid AS Process_Id, - R.tid AS Thread_Id, - R.stack_id AS Correlation_Id, - R.start AS Start_Timestamp, - R.end AS End_Timestamp - FROM "rocpd_info_node" AS N - INNER JOIN rocpd_info_process as P - ON P.guid = N.guid - AND P.nid = N.id - INNER JOIN regions AS R - ON R.guid = P.guid - AND R.nid = P.nid - AND R.pid = P.pid + 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 - R.category LIKE 'HIP_%' + category LIKE 'HIP_%' ORDER BY - R.start ASC, R.end DESC + start ASC, end DESC """ write_sql_query_to_csv(importData, query, output_path, "out_hip_api_trace") @@ -266,26 +234,19 @@ def write_hsa_api_csv( query = """ SELECT - R.guid AS Guid, - R.category AS Domain, - R.name AS Function, - R.pid AS Process_Id, - R.tid AS Thread_Id, - R.stack_id AS Correlation_Id, - R.start AS Start_Timestamp, - R.end AS End_Timestamp - FROM "rocpd_info_node" AS N - INNER JOIN rocpd_info_process as P - ON P.guid = N.guid - AND P.nid = N.id - INNER JOIN regions AS R - ON R.guid = P.guid - AND R.nid = P.nid - AND R.pid = P.pid + 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 - R.category LIKE 'HSA_%' + category LIKE 'HSA_%' ORDER BY - R.start ASC, R.end DESC + start ASC, end DESC """ write_sql_query_to_csv(importData, query, output_path, "out_hsa_api_trace") @@ -295,30 +256,23 @@ def write_marker_api_csv( query = """ SELECT - R.guid AS Guid, - R.category AS Domain, + guid AS Guid, + category AS Domain, CASE - WHEN json_extract(R.extdata, '$.message') IS NOT NULL - THEN json_extract(R.extdata, '$.message') - ELSE R.name + WHEN json_extract(extdata, '$.message') IS NOT NULL + THEN json_extract(extdata, '$.message') + ELSE name END AS Function, - R.pid AS Process_Id, - R.tid AS Thread_Id, - R.stack_id AS Correlation_Id, - R.start AS Start_Timestamp, - R.end AS End_Timestamp - FROM "rocpd_info_node" AS N - INNER JOIN rocpd_info_process as P - ON P.guid = N.guid - AND P.nid = N.id - INNER JOIN regions_and_samples AS R - ON R.guid = P.guid - AND R.nid = P.nid - AND R.pid = P.pid + 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 - R.category LIKE 'MARKER_%' + category LIKE 'MARKER_%' ORDER BY - R.start ASC, R.end DESC + start ASC, end DESC """ write_sql_query_to_csv(importData, query, output_path, "out_marker_api_trace") @@ -328,36 +282,29 @@ def write_counters_csv( query = """ SELECT - C.guid AS Pid, - C.stack_id AS Correlation_Id, - C.dispatch_id AS Dispatch_Id, - 'Agent ' || C.agent_log_index AS Agent_Id, - C.queue_id AS Queue_Id, - C.pid AS Process_Id, - C.tid AS Thread_Id, - C.grid_size AS Grid_Size, - C.kernel_id AS Kernel_Id, - C.kernel_name AS Kernel_Name, - C.workgroup_size AS Workgroup_Size, - C.lds_block_size AS LDS_Block_Size, - C.scratch_size AS Scratch_Size, - C.vgpr_count AS VGPR_Count, - C.accum_vgpr_count AS Accum_VGPR_Count, - C.sgpr_count AS SGPR_Count, - C.counter_name AS Counter_Name, - C.value AS Counter_Value, - C.start AS Start_Timestamp, - C.end AS End_Timestamp - FROM "rocpd_info_node" AS N - INNER JOIN rocpd_info_process as P - ON P.guid = N.guid - AND P.nid = N.id - INNER JOIN counters_collection AS C - ON C.guid = P.guid - AND C.nid = P.nid - AND C.pid = P.pid + guid AS Pid, + stack_id AS Correlation_Id, + dispatch_id AS Dispatch_Id, + 'Agent ' || agent_log_index 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, + vgpr_count AS 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 - C.start ASC, C.end DESC + start ASC, end DESC """ write_sql_query_to_csv(importData, query, output_path, "out_counter_collection") @@ -368,23 +315,16 @@ def write_scratch_memory_csv( query = """ SELECT 'SCRATCH_MEMORY' AS Kind, - 'SCRATCH_MEMORY_' || S.operation AS Operation, - 'Agent ' || S.agent_log_index AS Agent_Id, - S.queue_id AS Queue_Id, - S.tid AS Thread_Id, - S.alloc_flags AS Alloc_Flags, - S.start AS Start_Timestamp, - S.end AS End_Timestamp - FROM "rocpd_info_node" AS N - INNER JOIN rocpd_info_process as P - ON P.guid = N.guid - AND P.nid = N.id - INNER JOIN scratch_memory AS S - ON S.guid = P.guid - AND S.nid = P.nid - AND S.pid = P.pid + 'SCRATCH_MEMORY_' || operation AS Operation, + 'Agent ' || agent_log_index 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 - S.start ASC, S.end + start ASC, end DESC """ write_sql_query_to_csv(importData, query, output_path, "out_scratch_memory_trace") @@ -394,26 +334,19 @@ def write_rccl_api_csv( query = """ SELECT - R.guid AS Guid, - R.category AS Domain, - R.name AS Function, - R.pid AS Process_Id, - R.tid AS Thread_Id, - R.stack_id AS Correlation_Id, - R.start AS Start_Timestamp, - R.end AS End_Timestamp - FROM "rocpd_info_node" AS N - INNER JOIN rocpd_info_process as P - ON P.guid = N.guid - AND P.nid = N.id - INNER JOIN regions AS R - ON R.guid = P.guid - AND R.nid = P.nid - AND R.pid = P.pid + 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 - R.category LIKE 'RCCL_%' + category LIKE 'RCCL_%' ORDER BY - R.start ASC, R.end + start ASC, end DESC """ write_sql_query_to_csv(importData, query, output_path, "out_rccl_api_trace") @@ -423,26 +356,19 @@ def write_rocdecode_api_csv( query = """ SELECT - R.guid AS Guid, - R.category AS Domain, - R.name AS Function, - R.pid AS Process_Id, - R.tid AS Thread_Id, - R.stack_id AS Correlation_Id, - R.start AS Start_Timestamp, - R.end AS End_Timestamp - FROM "rocpd_info_node" AS N - INNER JOIN rocpd_info_process as P - ON P.guid = N.guid - AND P.nid = N.id - INNER JOIN regions AS R - ON R.guid = P.guid - AND R.nid = P.nid - AND R.pid = P.pid + 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 - R.category LIKE 'ROCDECODE_%' + category LIKE 'ROCDECODE_%' ORDER BY - R.start ASC, R.end + start ASC, end DESC """ write_sql_query_to_csv(importData, query, output_path, "out_rocdecode_api_trace") @@ -452,26 +378,19 @@ def write_rocjpeg_api_csv( query = """ SELECT - R.guid AS Guid, - R.category AS Domain, - R.name AS Function, - R.pid AS Process_Id, - R.tid AS Thread_Id, - R.stack_id AS Correlation_Id, - R.start AS Start_Timestamp, - R.end AS End_Timestamp - FROM "rocpd_info_node" AS N - INNER JOIN rocpd_info_process as P - ON P.guid = N.guid - AND P.nid = N.id - INNER JOIN regions AS R - ON R.guid = P.guid - AND R.nid = P.nid - AND R.pid = P.pid + 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 - R.category LIKE 'ROCJPEG_%' + category LIKE 'ROCJPEG_%' ORDER BY - R.start ASC, R.end + start ASC, end DESC """ write_sql_query_to_csv(importData, query, output_path, "out_rocjpeg_api_trace") From 56dfe6645966c7b68e1ed48929b4d503606c1511 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Thu, 31 Jul 2025 17:05:55 +0000 Subject: [PATCH 14/23] Format agent id --- source/lib/python/rocpd/csv.py | 143 ++++++++++++++++++++++----------- 1 file changed, 96 insertions(+), 47 deletions(-) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 7a1737fd34..7538c9bda4 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -41,7 +41,7 @@ def write_sql_query_to_csv( export_sqlite_query(connection, query, export_format="csv", export_path=export_path) def write_agent_info_csv( - importData, output_path + importData, config ) -> None: query = """ @@ -126,17 +126,26 @@ def write_agent_info_csv( model_name AS Model_Name FROM "rocpd_info_agent" """ - write_sql_query_to_csv(importData, query, output_path, "out_agent_info") + write_sql_query_to_csv(importData, query, config.output_path, "out_agent_info") def write_kernel_csv( - importData, output_path + importData, config ) -> None: - query = """ + 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, 'KERNEL_DISPATCH' AS Kind, - 'Agent ' || agent_log_index AS Agent_Id, + {agent_id} AS Agent_Id, queue_id AS Queue_Id, stream_id AS Stream_Id, tid AS Thread_Id, @@ -158,20 +167,33 @@ def write_kernel_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, output_path, "out_kernel_trace") + write_sql_query_to_csv(importData, query, config.output_path, "out_kernel_trace") def write_memory_copy_csv( - importData, output_path + importData, config ) -> None: - query = """ + 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, - 'Agent ' || src_agent_log_index AS Source_Agent_Id, - 'Agent ' || dst_agent_log_index AS Destination_Agent_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 @@ -179,20 +201,29 @@ def write_memory_copy_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, output_path, "out_memory_copy_trace") + write_sql_query_to_csv(importData, query, config.output_path, "out_memory_copy_trace") def write_memory_allocation_csv( - importData, output_path + importData, config ) -> None: - query = """ + 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 ' || agent_log_index + THEN {agent_id} ELSE '"' END AS Agent_Id, size AS Allocation_Size, @@ -204,10 +235,10 @@ def write_memory_allocation_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, output_path, "out_memory_allocation_trace") + write_sql_query_to_csv(importData, query, config.output_path, "out_memory_allocation_trace") def write_hip_api_csv( - importData, output_path + importData, config ) -> None: query = """ @@ -226,10 +257,10 @@ def write_hip_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, output_path, "out_hip_api_trace") + write_sql_query_to_csv(importData, query, config.output_path, "out_hip_api_trace") def write_hsa_api_csv( - importData, output_path + importData, config ) -> None: query = """ @@ -248,10 +279,10 @@ def write_hsa_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, output_path, "out_hsa_api_trace") + write_sql_query_to_csv(importData, query, config.output_path, "out_hsa_api_trace") def write_marker_api_csv( - importData, output_path + importData, config ) -> None: query = """ @@ -274,18 +305,27 @@ def write_marker_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, output_path, "out_marker_api_trace") + write_sql_query_to_csv(importData, query, config.output_path, "out_marker_api_trace") def write_counters_csv( - importData, output_path + importData, config ) -> None: - query = """ + 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 Pid, stack_id AS Correlation_Id, dispatch_id AS Dispatch_Id, - 'Agent ' || agent_log_index AS Agent_Id, + {agent_id} AS Agent_Id, queue_id AS Queue_Id, pid AS Process_Id, tid AS Thread_Id, @@ -306,17 +346,26 @@ def write_counters_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, output_path, "out_counter_collection") + write_sql_query_to_csv(importData, query, config.output_path, "out_counter_collection") def write_scratch_memory_csv( - importData, output_path + importData, config ) -> None: - query = """ + 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 'SCRATCH_MEMORY' AS Kind, 'SCRATCH_MEMORY_' || operation AS Operation, - 'Agent ' || agent_log_index AS Agent_Id, + {agent_id} AS Agent_Id, queue_id AS Queue_Id, tid AS Thread_Id, alloc_flags AS Alloc_Flags, @@ -326,10 +375,10 @@ def write_scratch_memory_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, output_path, "out_scratch_memory_trace") + write_sql_query_to_csv(importData, query, config.output_path, "out_scratch_memory_trace") def write_rccl_api_csv( - importData, output_path + importData, config ) -> None: query = """ @@ -348,10 +397,10 @@ def write_rccl_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, output_path, "out_rccl_api_trace") + write_sql_query_to_csv(importData, query, config.output_path, "out_rccl_api_trace") def write_rocdecode_api_csv( - importData, output_path + importData, config ) -> None: query = """ @@ -370,10 +419,10 @@ def write_rocdecode_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, output_path, "out_rocdecode_api_trace") + write_sql_query_to_csv(importData, query, config.output_path, "out_rocdecode_api_trace") def write_rocjpeg_api_csv( - importData, output_path + importData, config ) -> None: query = """ @@ -392,22 +441,22 @@ def write_rocjpeg_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, output_path, "out_rocjpeg_api_trace") + write_sql_query_to_csv(importData, query, config.output_path, "out_rocjpeg_api_trace") def write_csv(importData, config): - write_agent_info_csv(importData, config.output_path) - write_kernel_csv(importData, config.output_path) - write_memory_copy_csv(importData, config.output_path) - write_memory_allocation_csv(importData, config.output_path) - write_hip_api_csv(importData, config.output_path) - write_hsa_api_csv(importData, config.output_path) - write_marker_api_csv(importData, config.output_path) - write_counters_csv(importData, config.output_path) - write_scratch_memory_csv(importData, config.output_path) - write_rccl_api_csv(importData, config.output_path) - write_rocdecode_api_csv(importData, config.output_path) - write_rocjpeg_api_csv(importData, config.output_path) + 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): From cfc99557bdf323408f1b44bf7ac8d16779ed9b6b Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Thu, 31 Jul 2025 17:25:09 +0000 Subject: [PATCH 15/23] Compose full file name is sql writer function --- source/lib/python/rocpd/csv.py | 29 +++++++++++++++-------------- 1 file changed, 15 insertions(+), 14 deletions(-) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 7538c9bda4..1960d21466 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -32,12 +32,13 @@ from . import libpyrocpd def write_sql_query_to_csv( - connection: RocpdImportData, query, output_path, filename="" + 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.""" # call query module to export to csv - export_path = os.path.join(output_path, f"{filename}.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( @@ -126,7 +127,7 @@ def write_agent_info_csv( model_name AS Model_Name FROM "rocpd_info_agent" """ - write_sql_query_to_csv(importData, query, config.output_path, "out_agent_info") + write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "agent_info") def write_kernel_csv( importData, config @@ -167,7 +168,7 @@ def write_kernel_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, "out_kernel_trace") + write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "kernel") def write_memory_copy_csv( importData, config @@ -201,7 +202,7 @@ def write_memory_copy_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, "out_memory_copy_trace") + write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "memory_copy") def write_memory_allocation_csv( importData, config @@ -235,7 +236,7 @@ def write_memory_allocation_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, "out_memory_allocation_trace") + write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "memory_allocation") def write_hip_api_csv( importData, config @@ -257,7 +258,7 @@ def write_hip_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, "out_hip_api_trace") + write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "hip_api") def write_hsa_api_csv( importData, config @@ -279,7 +280,7 @@ def write_hsa_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, "out_hsa_api_trace") + write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "hsa_api") def write_marker_api_csv( importData, config @@ -305,7 +306,7 @@ def write_marker_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, "out_marker_api_trace") + write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "marker_api") def write_counters_csv( importData, config @@ -346,7 +347,7 @@ def write_counters_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, "out_counter_collection") + write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "counter_collection") def write_scratch_memory_csv( importData, config @@ -375,7 +376,7 @@ def write_scratch_memory_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, "out_scratch_memory_trace") + write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "scratch_memory") def write_rccl_api_csv( importData, config @@ -397,7 +398,7 @@ def write_rccl_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, "out_rccl_api_trace") + write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "rccl_api") def write_rocdecode_api_csv( importData, config @@ -419,7 +420,7 @@ def write_rocdecode_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, "out_rocdecode_api_trace") + write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "rocdecode_api") def write_rocjpeg_api_csv( importData, config @@ -441,7 +442,7 @@ def write_rocjpeg_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, "out_rocjpeg_api_trace") + write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "rocjpeg_api") def write_csv(importData, config): From 210780ffac21667957c445bff1b2e48932c52b10 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Thu, 31 Jul 2025 18:31:58 +0000 Subject: [PATCH 16/23] Add missing fields to kernel traces csv --- source/lib/python/rocpd/csv.py | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 1960d21466..4be7318a8d 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -156,8 +156,11 @@ def write_kernel_csv( stack_id AS Correlation_Id, start AS Start_Timestamp, end AS End_Timestamp, - scratch_size AS Private_Segment_Size, - lds_size AS Group_Segment_Size, + lds_size AS LDS_Block_Size, + scratch_size AS Scratch_Size, + vgpr_count AS 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, @@ -323,7 +326,7 @@ def write_counters_csv( query = f""" SELECT - guid AS Pid, + guid AS Guid, stack_id AS Correlation_Id, dispatch_id AS Dispatch_Id, {agent_id} AS Agent_Id, @@ -364,6 +367,7 @@ def write_scratch_memory_csv( query = f""" SELECT + guid AS Guid, 'SCRATCH_MEMORY' AS Kind, 'SCRATCH_MEMORY_' || operation AS Operation, {agent_id} AS Agent_Id, From 4fcad478f2fc12f2e4e6b955b0e151f329c2064c Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Mon, 4 Aug 2025 12:55:49 +0000 Subject: [PATCH 17/23] Rename vgpr_count to arch_vgpr_count --- source/lib/python/rocpd/csv.py | 4 ++-- source/lib/python/rocpd/source/csv.cpp | 4 ++-- source/lib/python/rocpd/source/types.hpp | 8 ++++---- source/share/rocprofiler-sdk-rocpd/data_views.sql | 4 ++-- 4 files changed, 10 insertions(+), 10 deletions(-) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 4be7318a8d..5a8be09017 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -158,7 +158,7 @@ def write_kernel_csv( end AS End_Timestamp, lds_size AS LDS_Block_Size, scratch_size AS Scratch_Size, - vgpr_count AS VGPR_Count, + 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, @@ -339,7 +339,7 @@ def write_counters_csv( workgroup_size AS Workgroup_Size, lds_block_size AS LDS_Block_Size, scratch_size AS Scratch_Size, - vgpr_count AS VGPR_Count, + arch_vgpr_count AS Arch_VGPR_Count, accum_vgpr_count AS Accum_VGPR_Count, sgpr_count AS SGPR_Count, counter_name AS Counter_Name, diff --git a/source/lib/python/rocpd/source/csv.cpp b/source/lib/python/rocpd/source/csv.cpp index 5cc5d0ac96..aa5b947d1a 100644 --- a/source/lib/python/rocpd/source/csv.cpp +++ b/source/lib/python/rocpd/source/csv.cpp @@ -258,7 +258,7 @@ write_kernel_csv( kernel.end, kernel.lds_size, kernel.scratch_size, - kernel.vgpr_count, + kernel.arch_vgpr_count, kernel.accum_vgpr_count, kernel.sgpr_count, kernel.workgroup_size.x, @@ -695,7 +695,7 @@ write_counters_csv(CsvManager& cs counter.workgroup_size, counter.lds_block_size, counter.scratch_size, - counter.vgpr_count, + counter.arch_vgpr_count, counter.accum_vgpr_count, counter.sgpr_count, fmt::format("\"{}\"", counter.counter_name), 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, From ea5d05b5bea486c92122dc83c55e6c83525f3c22 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Mon, 4 Aug 2025 13:09:59 +0000 Subject: [PATCH 18/23] Fix kernel name --- source/lib/python/rocpd/csv.py | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 5a8be09017..6bec8c01c4 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -142,6 +142,11 @@ def write_kernel_csv( else: agent_id = "" + if config.kernel_rename : + kernel_name = "region" + else: + kernel_name = "name" + query = f""" SELECT guid AS Guid, @@ -152,7 +157,7 @@ def write_kernel_csv( tid AS Thread_Id, dispatch_id AS Dispatch_Id, kernel_Id AS Kernel_Id, - name AS Kernel_Name, + {kernel_name} AS Kernel_Name, stack_id AS Correlation_Id, start AS Start_Timestamp, end AS End_Timestamp, From 58e02611217b9947ec0edf2568c94e359dabc873 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Mon, 4 Aug 2025 13:23:31 +0000 Subject: [PATCH 19/23] Skip empty query results --- source/lib/python/rocpd/csv.py | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 6bec8c01c4..5e0697f90d 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -36,6 +36,12 @@ def write_sql_query_to_csv( ) -> 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") From 32ccfb83712c9e2ee6b14688c80efc75c17713cb Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Mon, 4 Aug 2025 18:54:44 +0000 Subject: [PATCH 20/23] Format csv.py --- source/lib/python/rocpd/csv.py | 163 ++++++++++++++++++++------------- 1 file changed, 99 insertions(+), 64 deletions(-) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 5e0697f90d..942d0c6ad4 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -31,6 +31,7 @@ from . import output_config from . import libpyrocpd + def write_sql_query_to_csv( connection: RocpdImportData, query, output_path, output_file, filename="" ) -> None: @@ -47,9 +48,8 @@ def write_sql_query_to_csv( 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: + +def write_agent_info_csv(importData, config) -> None: query = """ SELECT @@ -133,22 +133,27 @@ def write_agent_info_csv( model_name AS Model_Name FROM "rocpd_info_agent" """ - write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "agent_info") + 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 +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) + 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 + 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 : + if config.kernel_rename: kernel_name = "region" else: kernel_name = "name" @@ -182,19 +187,24 @@ def write_kernel_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "kernel") + 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 +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) + 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 + 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: @@ -216,17 +226,22 @@ def write_memory_copy_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "memory_copy") + 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 +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) + 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 + elif ( + config.agent_index_value == libpyrocpd.agent_indexing.logical_node_type + ): # type-relative agent_id = "agent_type || ' ' || agent_type_index" else: agent_id = "" @@ -250,11 +265,12 @@ def write_memory_allocation_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "memory_allocation") + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "memory_allocation" + ) -def write_hip_api_csv( - importData, config -) -> None: + +def write_hip_api_csv(importData, config) -> None: query = """ SELECT @@ -272,11 +288,12 @@ def write_hip_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "hip_api") + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "hip_api" + ) -def write_hsa_api_csv( - importData, config -) -> None: + +def write_hsa_api_csv(importData, config) -> None: query = """ SELECT @@ -294,11 +311,12 @@ def write_hsa_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "hsa_api") + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "hsa_api" + ) -def write_marker_api_csv( - importData, config -) -> None: + +def write_marker_api_csv(importData, config) -> None: query = """ SELECT @@ -320,17 +338,22 @@ def write_marker_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "marker_api") + 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 +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) + 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 + elif ( + config.agent_index_value == libpyrocpd.agent_indexing.logical_node_type + ): # type-relative agent_id = "agent_type || ' ' || agent_type_index" else: agent_id = "" @@ -361,17 +384,22 @@ def write_counters_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "counter_collection") + 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 +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) + 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 + elif ( + config.agent_index_value == libpyrocpd.agent_indexing.logical_node_type + ): # type-relative agent_id = "agent_type || ' ' || agent_type_index" else: agent_id = "" @@ -391,11 +419,12 @@ def write_scratch_memory_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "scratch_memory") + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "scratch_memory" + ) -def write_rccl_api_csv( - importData, config -) -> None: + +def write_rccl_api_csv(importData, config) -> None: query = """ SELECT @@ -413,11 +442,12 @@ def write_rccl_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "rccl_api") + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "rccl_api" + ) -def write_rocdecode_api_csv( - importData, config -) -> None: + +def write_rocdecode_api_csv(importData, config) -> None: query = """ SELECT @@ -435,11 +465,12 @@ def write_rocdecode_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "rocdecode_api") + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "rocdecode_api" + ) -def write_rocjpeg_api_csv( - importData, config -) -> None: + +def write_rocjpeg_api_csv(importData, config) -> None: query = """ SELECT @@ -457,7 +488,10 @@ def write_rocjpeg_api_csv( ORDER BY start ASC, end DESC """ - write_sql_query_to_csv(importData, query, config.output_path, config.output_file, "rocjpeg_api") + write_sql_query_to_csv( + importData, query, config.output_path, config.output_file, "rocjpeg_api" + ) + def write_csv(importData, config): @@ -474,6 +508,7 @@ def write_csv(importData, config): write_rocdecode_api_csv(importData, config) write_rocjpeg_api_csv(importData, config) + def execute(input, config=None, window_args=None, **kwargs): importData = RocpdImportData(input) From 17830408639f5ca7ace6bfc260c79c8a784dff70 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Mon, 4 Aug 2025 19:19:22 +0000 Subject: [PATCH 21/23] Delete c++ CSV writer --- source/lib/python/rocpd/libpyrocpd.cpp | 107 --- source/lib/python/rocpd/source/CMakeLists.txt | 15 +- source/lib/python/rocpd/source/csv.cpp | 737 ------------------ source/lib/python/rocpd/source/csv.hpp | 120 --- 4 files changed, 3 insertions(+), 976 deletions(-) delete mode 100644 source/lib/python/rocpd/source/csv.cpp delete mode 100644 source/lib/python/rocpd/source/csv.hpp 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 aa5b947d1a..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.arch_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.arch_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 From 314ccae5f1f22f529c5d1c6d7aaea30864a6e4e1 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Tue, 5 Aug 2025 13:44:43 +0000 Subject: [PATCH 22/23] Add CSV header comparison test --- tests/pytest-packages/tests/rocprofv3.py | 187 +++++++++++++++++++++++ tests/rocprofv3/rocpd/CMakeLists.txt | 8 +- tests/rocprofv3/rocpd/conftest.py | 15 ++ tests/rocprofv3/rocpd/validate.py | 19 +++ 4 files changed, 228 insertions(+), 1 deletion(-) 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) From 8546caa8cd998477335fa4a58ca8ccfef18ab516 Mon Sep 17 00:00:00 2001 From: Aleksei Tumakaev Date: Tue, 5 Aug 2025 14:00:37 +0000 Subject: [PATCH 23/23] Fix comment spacing in csv.py --- source/lib/python/rocpd/csv.py | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/source/lib/python/rocpd/csv.py b/source/lib/python/rocpd/csv.py index 942d0c6ad4..f8ae2f9e43 100644 --- a/source/lib/python/rocpd/csv.py +++ b/source/lib/python/rocpd/csv.py @@ -148,7 +148,7 @@ def write_kernel_csv(importData, config) -> None: agent_id = "'Agent ' || agent_log_index" elif ( config.agent_index_value == libpyrocpd.agent_indexing.logical_node_type - ): # type-relative + ): # type-relative agent_id = "agent_type || ' ' || agent_type_index" else: agent_id = "" @@ -204,7 +204,7 @@ def write_memory_copy_csv(importData, config) -> None: dst_agent_id = "'Agent ' || dst_agent_log_index" elif ( config.agent_index_value == libpyrocpd.agent_indexing.logical_node_type - ): # type-relative + ): # type-relative src_agent_id = "src_agent_type || ' ' || src_agent_type_index" dst_agent_id = "dst_agent_type || ' ' || dst_agent_type_index" else: @@ -241,7 +241,7 @@ def write_memory_allocation_csv(importData, config) -> None: agent_id = "'Agent ' || agent_log_index" elif ( config.agent_index_value == libpyrocpd.agent_indexing.logical_node_type - ): # type-relative + ): # type-relative agent_id = "agent_type || ' ' || agent_type_index" else: agent_id = "" @@ -353,7 +353,7 @@ def write_counters_csv(importData, config) -> None: agent_id = "'Agent ' || agent_log_index" elif ( config.agent_index_value == libpyrocpd.agent_indexing.logical_node_type - ): # type-relative + ): # type-relative agent_id = "agent_type || ' ' || agent_type_index" else: agent_id = "" @@ -399,7 +399,7 @@ def write_scratch_memory_csv(importData, config) -> None: agent_id = "'Agent ' || agent_log_index" elif ( config.agent_index_value == libpyrocpd.agent_indexing.logical_node_type - ): # type-relative + ): # type-relative agent_id = "agent_type || ' ' || agent_type_index" else: agent_id = ""