[Autobackout][FuncReg]Revert of change: 6a9e9d9

 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;
    }

IGC/Compiler/CISACodeGen/ShaderCodeGen.cpp

@@ -83,7 +83,6 @@ 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"
@@ -1430,13 +1429,6 @@ void OptimizeIR(CodeGenContext* const pContext)
  mpm.add(createSROAPass());
  }
  }
-
- if (pContext->type == ShaderType::OPENCL_SHADER &&
- pContext->m_retryManager.IsFirstTry())
- {
- mpm.add(createGEPLoopStrengthReductionPass(IGC_IS_FLAG_ENABLED(allowLICM) &&
- pContext->m_retryManager.AllowLICM()));
- }
  }
 
  // Note:

IGC/Compiler/InitializePasses.h

@@ -66,7 +66,6 @@ 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,7 +20,6 @@ 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,7 +69,6 @@ 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,7 +123,6 @@ 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,7 +170,6 @@ 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