Add GEP Loop Strength Reduction pass

Adds new optimization pass reducing strength of GEP instructions in loop. Xe architecture
doesn't have native support for access from "base pointer plus offset". Codegen must
translate GEP instruction "getelementptr %constantPointer, %variableOffset" to
"(long long)pointer + offset*sizeof(*pointer)", generating additional mov/add/shl instructions.
If GEP is inside a loop and offset is incremented in constant steps, it is beneficial to change
GEP into induction variable "getelementptr %variablePointer, %constantOffset".

This change adds new pass translating code:

    int id = get_global_id(0);
    for (int i = 32; i < n_iters - 32; i += 32) {
        output[id + i] = output[id + i + 32] * output[id + i - 32];
    }

Into:

    int id = get_global_id(0);
    global float* outputm32 = output + id;
    for (int i = 32; i < n_iters - 32; i += 32, outputm32 += 32) {
        *(outputm32 + 32) = *(outputm32 + 64) * *outputm32;
    }

Author: pkwasnie-intel

IGC/Compiler/CISACodeGen/ShaderCodeGen.cpp

@@ -83,6 +83,7 @@ SPDX-License-Identifier: MIT
 #include "Compiler/Optimizer/OpenCLPasses/UnreachableHandling/UnreachableHandling.hpp"
 #include "Compiler/Optimizer/OpenCLPasses/WIFuncs/WIFuncResolution.hpp"
 #include "Compiler/Optimizer/OpenCLPasses/ScalarArgAsPointer/ScalarArgAsPointer.hpp"
+#include "Compiler/Optimizer/OpenCLPasses/GEPLoopStrengthReduction/GEPLoopStrengthReduction.hpp"
 #include "Compiler/Optimizer/MCSOptimization.hpp"
 #include "Compiler/Optimizer/GatingSimilarSamples.hpp"
 #include "Compiler/Optimizer/IntDivConstantReduction.hpp"
@@ -355,6 +356,15 @@ static void UpdateInstTypeHint(CodeGenContext& ctx)
 // forward declaration
 llvm::ModulePass* createPruneUnusedArgumentsPass();
 
+static bool useStatelessToStateful(CodeGenContext& ctx)
+{
+    return (ctx.m_instrTypes.hasLoadStore &&
+        ctx.m_DriverInfo.SupportsStatelessToStatefulBufferTransformation() &&
+        !ctx.getModuleMetaData()->compOpt.GreaterThan4GBBufferRequired &&
+        IGC_IS_FLAG_ENABLED(EnableStatelessToStateful) &&
+        !ctx.m_instrTypes.hasInlineAsmPointerAccess);
+}
+
 void AddLegalizationPasses(CodeGenContext& ctx, IGCPassManager& mpm, PSSignature* pSignature)
 {
     COMPILER_TIME_START(&ctx, TIME_CG_Add_Legalization_Passes);
@@ -705,12 +715,7 @@ void AddLegalizationPasses(CodeGenContext& ctx, IGCPassManager& mpm, PSSignature
         mpm.add(new PromoteStatelessToBindless());
     }
 
-    if (!isOptDisabled &&
-        ctx.m_instrTypes.hasLoadStore &&
-        ctx.m_DriverInfo.SupportsStatelessToStatefulBufferTransformation() &&
-        !ctx.getModuleMetaData()->compOpt.GreaterThan4GBBufferRequired &&
-        IGC_IS_FLAG_ENABLED(EnableStatelessToStateful) &&
-        !ctx.m_instrTypes.hasInlineAsmPointerAccess)
+    if (!isOptDisabled && useStatelessToStateful(ctx))
     {
         mpm.add(new StatelessToStateful());
     }
@@ -1485,6 +1490,15 @@ void OptimizeIR(CodeGenContext* const pContext)
                         mpm.add(createSROAPass());
                     }
                 }
+
+                if (pContext->type == ShaderType::OPENCL_SHADER &&
+                    pContext->platform.isCoreChildOf(IGFX_XE_HPC_CORE) &&
+                    !useStatelessToStateful(*pContext) &&
+                    pContext->m_retryManager.IsFirstTry())
+                {
+                    mpm.add(createGEPLoopStrengthReductionPass(IGC_IS_FLAG_ENABLED(allowLICM) &&
+                            pContext->m_retryManager.AllowLICM()));
+                }
             }
 
             // Note:

IGC/Compiler/InitializePasses.h

@@ -66,6 +66,7 @@ void initializeGenericAddressAnalysisPass(llvm::PassRegistry&);
 void initializeGenericAddressDynamicResolutionPass(llvm::PassRegistry&);
 void initializeGenFDIVEmulationPass(llvm::PassRegistry&);
 void initializeGenIRLoweringPass(llvm::PassRegistry&);
+void initializeGEPLoopStrengthReductionPass(llvm::PassRegistry&);
 void initializeGEPLoweringPass(llvm::PassRegistry&);
 void initializeGenSpecificPatternPass(llvm::PassRegistry&);
 void initializeGreedyLiveRangeReductionPass(llvm::PassRegistry&);

IGC/Compiler/Optimizer/OpenCLPasses/CMakeLists.txt

@@ -20,6 +20,7 @@ add_subdirectory(DeviceEnqueueFuncs)
 add_subdirectory(DisableLoopUnrollOnRetry)
 add_subdirectory(ExtensionFuncs)
 add_subdirectory(GenericAddressResolution)
+add_subdirectory(GEPLoopStrengthReduction)
 add_subdirectory(ImageFuncs)
 add_subdirectory(LocalBuffers)
 add_subdirectory(NamedBarriers)
@@ -70,6 +71,7 @@ set(IGC_BUILD__SRC__Optimizer_OpenCLPasses_All
     ${IGC_BUILD__SRC__OpenCLPasses_DpasFuncs}
     ${IGC_BUILD__SRC__OpenCLPasses_ExtensionFuncs}
     ${IGC_BUILD__SRC__OpenCLPasses_GenericAddressResolution}
+    ${IGC_BUILD__SRC__OpenCLPasses_GEPLoopStrengthReduction}
     ${IGC_BUILD__SRC__OpenCLPasses_ImageFuncs}
     ${IGC_BUILD__SRC__OpenCLPasses_LocalBuffers}
     ${IGC_BUILD__SRC__OpenCLPasses_LowerLocalMemPool}
@@ -125,6 +127,7 @@ set(IGC_BUILD__HDR__Optimizer_OpenCLPasses_All
     ${IGC_BUILD__HDR__OpenCLPasses_DpasFuncs}
     ${IGC_BUILD__HDR__OpenCLPasses_ExtensionFuncs}
     ${IGC_BUILD__HDR__OpenCLPasses_GenericAddressResolution}
+    ${IGC_BUILD__HDR__OpenCLPasses_GEPLoopStrengthReduction}
     ${IGC_BUILD__HDR__OpenCLPasses_ImageFuncs}
     ${IGC_BUILD__HDR__OpenCLPasses_LocalBuffers}
     ${IGC_BUILD__HDR__OpenCLPasses_LowerLocalMemPool}
@@ -173,6 +176,7 @@ set(IGC_BUILD_Compiler_OpenCLPasses_Groups
     Compiler__OpenCLPasses_ExtensionFuncs
     Compiler__OpenCLPasses_GenericAddressResolution
     Compiler__OpenCLPasses_GenericAddressSpaceStaticResolution
+    Compiler__OpenCLPasses_GEPLoopStrengthReduction
     Compiler__OpenCLPasses_ImageFuncs
     Compiler__OpenCLPasses_LocalBuffers
     Compiler__OpenCLPasses_LowerLocalMemPool

IGC/Compiler/Optimizer/OpenCLPasses/GEPLoopStrengthReduction/CMakeLists.txt

@@ -0,0 +1,29 @@
+#=========================== begin_copyright_notice ============================
+#
+# Copyright (C) 2023 Intel Corporation
+#
+# SPDX-License-Identifier: MIT
+#
+#============================ end_copyright_notice =============================
+
+include_directories("${CMAKE_CURRENT_SOURCE_DIR}")
+
+
+set(IGC_BUILD__SRC__GEPLoopStrengthReduction
+    "${CMAKE_CURRENT_SOURCE_DIR}/GEPLoopStrengthReduction.cpp"
+  )
+set(IGC_BUILD__SRC__OpenCLPasses_GEPLoopStrengthReduction ${IGC_BUILD__SRC__GEPLoopStrengthReduction} PARENT_SCOPE)
+
+set(IGC_BUILD__HDR__GEPLoopStrengthReduction
+    "${CMAKE_CURRENT_SOURCE_DIR}/GEPLoopStrengthReduction.hpp"
+  )
+set(IGC_BUILD__HDR__OpenCLPasses_GEPLoopStrengthReduction ${IGC_BUILD__HDR__GEPLoopStrengthReduction} PARENT_SCOPE)
+
+
+igc_sg_register(
+    Compiler__OpenCLPasses_GEPLoopStrengthReduction
+    "GEPLoopStrengthReduction"
+    FILES
+      ${IGC_BUILD__SRC__GEPLoopStrengthReduction}
+      ${IGC_BUILD__HDR__GEPLoopStrengthReduction}
+  )