Add an OCL option to enable profile-guided trimming

Add the OCL option -ze-opt-profile-guided-trimming to enable
profile-guided trimming.

Author: tmpark

IGC/AdaptorOCL/UnifyIROCL.cpp

@@ -425,8 +425,10 @@ static void CommonOCLBasedPasses(OpenCLProgramContext* pContext)
         mpm.add(createSCCPPass());
         mpm.add(new ResolveConstExprCalls());
 
-        // Estimate maximal function size in the module and disable subroutine if not profitable.
-        mpm.add(createEstimateFunctionSizePass());
+        // Estimate maximal function size in the module and disable subroutine
+        // if not profitable.
+        mpm.add(createEstimateFunctionSizePass(
+            pContext->m_Options.StaticProfileGuidedTrimming));
         mpm.add(createProcessFuncAttributesPass());
         FastMathFlags Mask;
         Mask.setFast();

IGC/Compiler/CISACodeGen/EstimateFunctionSize.cpp

@@ -42,7 +42,7 @@ namespace IGC {
             AL_Kernel
         };
 
-        explicit EstimateFunctionSize(AnalysisLevel = AL_Module);
+        explicit EstimateFunctionSize(AnalysisLevel = AL_Module, bool = false);
         ~EstimateFunctionSize();
         EstimateFunctionSize(const EstimateFunctionSize&) = delete;
         EstimateFunctionSize& operator=(const EstimateFunctionSize&) = delete;
@@ -137,9 +137,45 @@ namespace IGC {
         llvm::ScaledNumber<uint64_t> thresholdForTrimming;
         std::unordered_map<llvm::Loop *, llvm::ScaledNumber<uint64_t>>
             LoopIterCnts;
+
+        // Flags for Kernel trimming
+        bool ControlKernelTotalSize;
+        bool ControlUnitSize;
+        unsigned ControlInlineTinySize;
+        unsigned UnitSizeThreshold;
+
+        // Flags for Static Profile-guided trimming
+        bool StaticProfileGuidedTrimming;
+        bool UseFrequencyInfoForSPGT;
+        bool BlockFrequencySampling;
+        bool EnableLeafCollapsing;
+        bool EnableSizeContributionOptimization;
+        bool LoopCountAwareTrimming;
+        bool EnableGreedyTrimming;
+        unsigned SizeWeightForSPGT;
+        unsigned FrequencyWeightForSPGT;
+        unsigned MetricForKernelSizeReduction;
+        unsigned ParameterForColdFuncThreshold;
+        unsigned ControlInlineTinySizeForSPGT;
+        unsigned MaxUnrollCountForFunctionSizeAnalysis;
+        unsigned SkipTrimmingOneCopyFunction;
+        std::string SelectiveTrimming;
+
+        // Flags for Partitioning
+        bool PartitionUnit;
+        bool StaticProfileGuidedPartitioning;
+
+        // Flags for implcit arguments and external functions
+        bool ForceInlineExternalFunctions;
+        bool ForceInlineStackCallWithImplArg;
+        bool ControlInlineImplicitArgs;
+        unsigned SubroutineThreshold;
+        unsigned KernelTotalSizeThreshold;
+        unsigned ExpandedUnitSizeThreshold;
     };
 
     llvm::ModulePass* createEstimateFunctionSizePass();
+    llvm::ModulePass *createEstimateFunctionSizePass(bool);
     llvm::ModulePass* createEstimateFunctionSizePass(EstimateFunctionSize::AnalysisLevel);
 
 } // namespace IGC

IGC/Compiler/CISACodeGen/EstimateFunctionSize.h

@@ -603,5 +603,9 @@ void Options::parseOptions(const char* opts)
         Xfinalizer = true;
         XfinalizerOption = arg->getValue();
     }
+
+    if (apiOptions.hasArg(OPT_static_profile_guided_trimming_common)) {
+        StaticProfileGuidedTrimming = true;
+    }
 }
 } // namespace IGC

IGC/Compiler/CISACodeGen/OpenCLOptions.cpp

@@ -191,6 +191,8 @@ class Options
     // This option enables FP64 emulation for conversions
     // This applies to platforms that cannot HW support for double operations
     bool EnableFP64GenConvEmu = false;
+    // This option enables static profile-guided trimming
+    bool StaticProfileGuidedTrimming = false;
 
 private:
     void parseOptions(const char* opts);

IGC/Compiler/CISACodeGen/OpenCLOptions.hpp

@@ -1374,7 +1374,9 @@ namespace TC
                             (strcmp(pParam, "-ze-fp64-gen-emu") == 0) || //used by fp64 emulation
                             (strcmp(pParam, "-cl-fp64-gen-conv-emu") == 0) || //used by fp64 conversion emulation
                             (strcmp(pParam, "-ze-fp64-gen-conv-emu") == 0) || //used by fp64 conversion emulation
-                            (strcmp(pParam, "-Xfinalizer") == 0); // used to pass options to visa finalizer
+                            (strcmp(pParam, "-Xfinalizer") == 0) || // used to pass options to visa finalizer
+                            (strcmp(pParam, "-cl-intel-static-profile-guided-trimming") == 0) || //used to enable profile-guided trimming
+                            (strcmp(pParam, "-ze-opt-static-profile-guided-trimming") == 0); //used to enable profile-guided trimming
 
                         if (isCommonOption)
                         {

IGC/OCLFE/igd_fcl_mcl/source/clang_tb.cpp

@@ -58,6 +58,9 @@ defm fp64_gen_emu : CommonFlag<"fp64-gen-emu">;
 // -cl-fp64-gen-conv-emu  -ze-fp64-gen-conv-emu
 defm fp64_gen_conv_emu : CommonFlag<"fp64-gen-conv-emu">;
 
+// -cl-intel-profile-guided-trimming, -ze-opt-profile-guided-trimming
+defm static_profile_guided_trimming : CommonFlag<"static-profile-guided-trimming">;
+
 // }} Backend API options
 
 // API options from source translation {{

IGC/Options/include/igc/Options/IGCApiOptions.td

@@ -0,0 +1,65 @@
+/*========================== begin_copyright_notice ============================
+
+Copyright (C) 2024 Intel Corporation
+
+SPDX-License-Identifier: MIT
+
+============================= end_copyright_notice ===========================*/
+
+// REQUIRES: regkeys, dg2-supported
+
+// RUN: ocloc compile -file %s -device dg2 -options "-igc_opts 'SubroutineThreshold=1,KernelTotalSizeThreshold=1,ControlInlineTinySize=1,PrintControlKernelTotalSize=15'" 2>&1 | FileCheck %s --check-prefix=CHECK-DEFAULT
+// RUN: ocloc compile -file %s -device dg2 -options "-ze-opt-static-profile-guided-trimming -igc_opts 'SubroutineThreshold=1,KernelTotalSizeThreshold=1,ControlInlineTinySize=1,ControlInlineTinySizeForSPGT=150,PrintControlKernelTotalSize=15'" 2>&1 | FileCheck %s --check-prefix=CHECK-SPGT
+
+// CHECK-DEFAULT: Good to trim (Big enough > 1), bar3, Function Attribute: Best effort innline, Function size: 13, Freq: 0.0
+// CHECK-SPGT: Can't trim (Low weight < 0.03218650818), bar3, Function Attribute: Best effort innline, Function size: 163, Size after collapsing: 163, Size contribution: 652, Freq: 327680.0, Weight: 0.002581326365
+
+int bar3(__global int *c) {
+    int k = 10;
+    for (int i = 0 ; i < 100 ; i++) {
+        *c += k * i;
+    }
+    return k;
+}
+
+int bar2(__global int *b) {
+    int k = 10;
+    for (int i = 0 ; i < 100 ; i++) {
+        *b += k * bar3(b);
+    }
+    return k;
+}
+
+int bar1(__global int *a) {
+    int k = 10;
+    for (int i = 0 ; i < 100 ; i++) {
+        *a += k * bar2(a);
+        *a += k * bar3(a);
+    }
+    return k;
+}
+
+__kernel void foo(int __global *p) {
+    int a = 0;
+    for (int i = 0; i < 100; i++) {
+        a += bar1(p);
+        a += bar2(p);
+        a += bar3(p);
+    }
+    for (int i = 300; i < 500000; i++) {
+        a += *p;
+    }
+    for (int i = 300; i < 500000; i++) {
+        a += *p;
+    }
+    for (int i = 300; i < 500000; i++) {
+        a += *p;
+    }
+    for (int i = 300; i < 500000; i++) {
+        a += *p;
+    }
+    for (int i = 300; i < 500000; i++) {
+        a += *p;
+    }
+    *p = a;
+}