enable IndVarSimplification pass

Enable IndVarSimplification pass.
intel · Sep 20, 2024 · 87d6432 · 87d6432
1 parent 0add600
commit 87d6432
Show file tree

Hide file tree

Showing 4 changed files with 53 additions and 0 deletions.
diff --git a/IGC/AdaptorOCL/dllInterfaceCompute.cpp b/IGC/AdaptorOCL/dllInterfaceCompute.cpp
@@ -1219,6 +1219,26 @@ bool TranslateBuildSPMD(
  }
  }
 
+ // From pass IndVarSimplify we are only interested in optimization done by -replexitval.
+ // Disable other features that can have a negative impact on performance.
+ std::array<llvm::StringRef, 4> indVarSimplifyFlags = {
+ "-indvars-post-increment-ranges=0",
+ "-disable-lftr=1",
+ "-indvars-widen-indvars=0",
+ "-verify-indvars=0"
+ };
+ for (auto indVarSimplifyFlag : indVarSimplifyFlags)
+ {
+ auto indVarSimplifySwitch = optionsMap.find(indVarSimplifyFlag.drop_front(1).split("=").first);
+ if (indVarSimplifySwitch != optionsMap.end())
+ {
+ if (indVarSimplifySwitch->getValue()->getNumOccurrences() == 0)
+ {
+ args.push_back(indVarSimplifyFlag.data());
+ }
+ }
+ }
+
  if (std::size(args) > 1)
  {
  llvm::cl::ParseCommandLineOptions(std::size(args), &args[0]);

diff --git a/IGC/Compiler/CISACodeGen/ShaderCodeGen.cpp b/IGC/Compiler/CISACodeGen/ShaderCodeGen.cpp
@@ -1525,6 +1525,11 @@ void OptimizeIR(CodeGenContext* const pContext)
 
  mpm.add(createIGCInstructionCombiningPass());
 
+ if (IGC_IS_FLAG_ENABLED(EnableIndVarSimplification))
+ {
+ mpm.add(llvm::createIndVarSimplifyPass());
+ }
+
  if (IGC_IS_FLAG_ENABLED(EnableLoopHoistConstant))
  {
  mpm.add(createLoopHoistConstant());

diff --git a/IGC/common/igc_flags.h b/IGC/common/igc_flags.h
@@ -209,6 +209,7 @@ DECLARE_IGC_REGKEY(bool, DisableIRVerification, false, "Setting this to
 DECLARE_IGC_REGKEY(bool, EnableJumpThreading, true, "Setting this to 1/true adds a compiler switch to enable llvm jumpThreading pass.", true)
 DECLARE_IGC_REGKEY(bool, DisableLoopUnroll, false, "Setting this to 1/true adds a compiler switch to disable loop unrolling.", true)
 DECLARE_IGC_REGKEY(DWORD, RuntimeLoopUnrolling, 0, "Setting this to switch on/off runtime loop unrolling. 0: default (on), 1: force on, 2: force off", false)
+DECLARE_IGC_REGKEY(bool, EnableIndVarSimplification, true, "Enables IndVarSimplification pass.", true)
 DECLARE_IGC_REGKEY(bool, DisableBranchSwaping, false, "Setting this to 1/true adds a compiler switch to disable branch swapping.", false)
 DECLARE_IGC_REGKEY(bool, DisableSynchronizationObjectCoalescingPass, false, "Disable SynchronizationObjectCoalescing pass", false)
 DECLARE_IGC_REGKEY(bool, EnableIndependentSharedMemoryFenceFunctionality, false, "Enable treating global memory fences as shared memory fences in SynchronizationObjectCoalescing pass", false)

diff --git a/IGC/ocloc_tests/optimizations/IndVarSimplification.cl b/IGC/ocloc_tests/optimizations/IndVarSimplification.cl
@@ -0,0 +1,27 @@
+/*========================== begin_copyright_notice ============================
+
+Copyright (C) 2024 Intel Corporation
+
+SPDX-License-Identifier: MIT
+
+============================= end_copyright_notice ===========================*/
+
+// REQUIRES: regkeys,pvc-supported,llvm-14-plus
+
+// RUN: ocloc compile -file %s -device pvc -options "-igc_opts 'PrintToConsole=1,PrintAfter=EmitPass'" 2>&1 | FileCheck %s
+
+// CHECK-NOT: phi
+// CHECK: [[LOAD:%.*]] = load i32, i32 addrspace(1)* [[PTR:%.*]]
+// CHECK: [[ADD:%.*]] = add i32 [[LOAD]], 10000
+// CHECK: store i32 [[ADD]], i32 addrspace(1)* [[PTR]]
+
+kernel void test(global int* ptr)
+{
+ size_t id = get_global_id(0);
+ int val = ptr[id];
+ for (int i = 0; i < 10000; i++)
+ {
+ val += 1;
+ }
+ ptr[id] = val;
+}