Enabling the TimeStatsPerPass option

Enabling the TimeStatsPerPass option

Author: ppogotov

IGC/AdaptorOCL/dllInterfaceCompute.cpp

@@ -1704,6 +1704,7 @@ bool TranslateBuildSPMD(
 
     COMPILER_TIME_END(&oclContext, TIME_TOTAL);
 
+    COMPILER_TIME_PER_PASS_PRINT(&oclContext, ShaderType::OPENCL_SHADER, oclContext.hash);
     COMPILER_TIME_PRINT(&oclContext, ShaderType::OPENCL_SHADER, oclContext.hash);
 
     COMPILER_TIME_DEL(&oclContext, m_compilerTimeStats);

IGC/Compiler/CISACodeGen/MemOpt2.cpp

@@ -434,7 +434,7 @@ char MemOpt2::ID = 0;
 FunctionPass *createMemOpt2Pass(int MLT) { return new MemOpt2(MLT); }
 
 #define PASS_FLAG "igc-memopt2"
-#define PASS_DESC "IGC Memory Optimization, the 2nd"
+#define PASS_DESC "IGC Memory Optimization the 2nd"
 #define PASS_CFG_ONLY false
 #define PASS_ANALYSIS false
 IGC_INITIALIZE_PASS_BEGIN(MemOpt2, PASS_FLAG, PASS_DESC, PASS_CFG_ONLY,

IGC/common/Stats.cpp

@@ -531,8 +531,6 @@ void TimeStats::printTime(ShaderType type, ShaderHash hash, void* context, UINT6
 
     pp.printTimeCSV( shaderName, psoDDIHash);
 
-    // Skip printing PerPass info to CSV for now
-    //pp.printPerPassTimeCSV( shaderName );
     // Avoid printing the timestats for each shader to screen
     // pp.printSumTimeTable(llvm::dbgs());
 }
@@ -572,6 +570,19 @@ void TimeStats::printSumTime() const
     pp.printSumTimeTable(llvm::dbgs());
 }
 
+void TimeStats::printTimePerPass(ShaderType type, ShaderHash hash) const
+{
+    TimeStats pp = postProcess();
+
+    std::string shaderName = IGC::Debug::DumpName(IGC::Debug::GetShaderOutputName()).Type(type).Hash(hash).StagedInfo(nullptr).str().c_str();
+    if (shaderName.find_last_of("\\") != std::string::npos)
+    {
+        shaderName = shaderName.substr(shaderName.find_last_of("\\") + 1, shaderName.size());
+    }
+
+    pp.printPerPassTimeCSV(shaderName);
+}
+
 bool TimeStats::skipTimer( int i ) const
 {
     const COMPILE_TIME_INTERVALS interval = static_cast<COMPILE_TIME_INTERVALS>(i);

IGC/common/Stats.hpp

@@ -218,6 +218,8 @@ class TimeStats
     void printTime( ShaderType type, ShaderHash hash, UINT64 psoDDIHash) const;
     void printTime( ShaderType type, ShaderHash hash, void* context, UINT64 psoDDIHash) const;
 
+    void printTimePerPass( ShaderType type, ShaderHash hash ) const;
+
     /// Print the aggregated times for multiple shaders
     void printSumTime() const;
     /// Print the times for all passes
@@ -426,6 +428,19 @@ class TimeStats
         } \
     } while (0)
 
+#define COMPILER_TIME_PER_PASS_PRINT( pointer, ...) \
+    do \
+    { \
+        if( (pointer) && (pointer)->m_compilerTimeStats ) \
+        { \
+            if ( IGC_REGKEY_OR_FLAG_ENABLED(DumpTimeStatsPerPass, TIME_STATS_PER_PASS ) ) \
+            { \
+                (pointer)->m_compilerTimeStats->printTimePerPass( \
+                    __VA_ARGS__ ); \
+            } \
+        } \
+    } while (0)
+
 #else // GET_TIME_STATS
 
 #   define COMPILER_TIME_START( pointer, value ) do { } while (0)

IGC/ocloc_tests/features/compiler_time_stats/time_stats_per_pass.cl

@@ -0,0 +1,16 @@
+/*========================== begin_copyright_notice ============================
+
+Copyright (C) 2024 Intel Corporation
+
+SPDX-License-Identifier: MIT
+
+============================= end_copyright_notice ===========================*/
+
+// REQUIRES: regkeys, pvc-supported
+
+// RUN: ocloc compile -file %s -options "-igc_opts 'DumpToCurrentDir=1 DumpTimeStatsPerPass=1'" -device pvc
+// RUN: FileCheck %s --input-file=TimeStatPerPass_Shaders.csv --check-prefixes=CHECK
+// CHECK: Frequency
+
+// This test checks that DumpTimeStatsPerPass option works correctly.
+__kernel void foo(int a, int b, __global int *res) { *res = a + b; }