Merge commit '7251cea0bc9e90217aea53bd8870a5be728814c2' into develop

projects/rocprofiler-sdk/CHANGELOG.md

@@ -218,6 +218,7 @@ Full documentation for ROCprofiler-SDK is available at [rocm.docs.amd.com/projec
 - Fixed queue preemption error and HSA_STATUS_ERROR_INVALID_PACKET_FORMAT error for stochastic PC-sampling for MI300X, leading to more stable runs.
 - Fixed the system hang issue for host-trap PC-sampling on MI300X. 
 - Fixed rocpd counter collection issue when counter collection alone is enabled, rocpd_kernel_dispatch table gets populated by counters data instead of kernel_dispatch data.
+- Fixed kernel trace csv output generated by rocpd.
 
 ### Removed
 

projects/rocprofiler-sdk/samples/thread_trace/CMakeLists.txt

@@ -21,6 +21,16 @@ project(rocprofiler-sdk-samples-thread-trace LANGUAGES CXX HIP)
 
 find_package(rocprofiler-sdk REQUIRED)
 find_package(rocprofiler-sdk-roctx REQUIRED)
+find_package(
+    amd_comgr
+    REQUIRED
+    CONFIG
+    HINTS
+    ${rocm_version_DIR}
+    ${ROCM_PATH}
+    PATHS
+    ${rocm_version_DIR}
+    ${ROCM_PATH})
 
 set(CMAKE_CXX_STANDARD 17)
 set(CMAKE_CXX_EXTENSIONS OFF)
@@ -44,11 +54,8 @@ target_sources(thread-trace-sample PRIVATE main.cpp agent.cpp)
 target_link_libraries(
     thread-trace-sample
     PUBLIC rocprofiler-sdk::samples-build-flags
-    PRIVATE rocprofiler-sdk::rocprofiler-sdk
-            rocprofiler-sdk::samples-common-library
-            rocprofiler-sdk::rocprofiler-sdk-dw
-            rocprofiler-sdk-roctx::rocprofiler-sdk-roctx
-            rocprofiler-sdk::rocprofiler-sdk-amd-comgr)
+    PRIVATE rocprofiler-sdk::rocprofiler-sdk rocprofiler-sdk::samples-common-library
+            rocprofiler-sdk-roctx::rocprofiler-sdk-roctx amd_comgr)
 
 add_test(NAME thread-trace-sample COMMAND $<TARGET_FILE:thread-trace-sample>)
 

projects/rocprofiler-sdk/source/lib/output/generateRocpd.cpp

@@ -1021,8 +1021,7 @@ write_rocpd(
             dispatch_evt_ids.at(dispatch_id) = evt_id;
 
             auto region_name =
-                (corr_id.external.value > 0 &&
-                 (enable_duplicate_check || kernel_id > 0))
+                (corr_id.external.value > 0 && (enable_duplicate_check || kernel_id > 0))
                     ? tool_metadata.get_kernel_name(kernel_id, corr_id.external.value)
                     : std::string_view{};
 
@@ -1088,7 +1087,7 @@ write_rocpd(
                                      dispatch_data.end_timestamp,      // end_timestamp
                                      info.grid_size,                   // grid
                                      info.workgroup_size,              // workgroup
-                                     false  // enable_duplicate_check
+                                     false                             // enable_duplicate_check
                     );
                 }
             }

projects/rocprofiler-sdk/source/lib/python/rocpd/source/csv.cpp

@@ -75,15 +75,11 @@ CsvManager::CsvManager(rocprofiler::tool::output_config output_cfg)
     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\",\"Private_Segment_Size\",\"Group_"
-          "Segment_Size\",\"Workgroup_Size_X\","
-          "\"Workgroup_Size_Y\",\"Workgroup_Size_Z\",\"Grid_"
-          "Size_X\",\"Grid_Size_Y\",\"Grid_Size_"
-          "Z\""}},
+          "\"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\","
@@ -260,8 +256,11 @@ write_kernel_csv(
                           kernel.stack_id,
                           kernel.start,
                           kernel.end,
-                          kernel.scratch_size,
                           kernel.lds_size,
+                          kernel.scratch_size,
+                          kernel.vgpr_count,
+                          kernel.accum_vgpr_count,
+                          kernel.sgpr_count,
                           kernel.workgroup_size.x,
                           kernel.workgroup_size.y,
                           kernel.workgroup_size.z,

projects/rocprofiler-sdk/source/lib/python/rocpd/source/types.hpp

@@ -286,6 +286,9 @@ 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                accum_vgpr_count    = 0;
+    uint64_t                sgpr_count          = 0;
 };
 
 struct memory_allocation
@@ -709,6 +712,9 @@ 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(accum_vgpr_count);
+    LOAD_DATA_FIELD(sgpr_count);
     LOAD_DATA_FIELD(static_lds_size);
     LOAD_DATA_FIELD(static_scratch_size);
     LOAD_DATA_FIELD(stack_id);

projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/data_views.sql

@@ -315,6 +315,9 @@ 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.accum_vgpr_count,
+    S.sgpr_count,
     S.group_segment_size AS static_lds_size,
     S.private_segment_size AS static_scratch_size,
     E.stack_id,