[Polly] Require LoadInst base ptrs to be hoisted.

[Meinersbur]

Summary

Only when load-hoisted we can be sure the base pointer is invariant during the SCoP's execution. Most of the time it would be added to the required hoists for the alias checks anyway, except with -polly-ignore-aliasing, -polly-use-runtime-alias-checks=0 or if AliasAnalysis is already sure it doesn't alias with anything.

Two more parts in Polly assume that this load-hoisting took place:

• setNewAccessRelation() which contains an assert which tests this.
• BlockGenerator which would use to the base ptr from the original code if not load-hoisted (if the access expression is regenerated)

polly/trunk/include/polly/ScopDetection.h

private:
   /// @return True if all loads can be assumed invariant.
   bool onlyValidRequiredInvariantLoads(InvariantLoadsSetTy &RequiredILS,
                                        DetectionContext &Context) const;
 
   /// Check if a value is invariant in the region Reg.
   ///
   /// @param Val Value to check for invariance.
   /// @param Reg The region to consider for the invariance of Val.
+  /// @param Ctx The current detection context.
   ///
   /// @return True if the value represented by Val is invariant in the region
   ///         identified by Reg.
-  bool isInvariant(const Value &Val, const Region &Reg) const;
+  bool isInvariant(Value &Val, const Region &Reg, DetectionContext &Ctx) const;
 
   /// Check if the memory access caused by @p Inst is valid.
   ///
   /// @param Inst    The access instruction.
   /// @param AF      The access function.
   /// @param BP      The access base pointer.
   /// @param Context The current detection context.
   bool isValidAccess(Instruction *Inst, const SCEV *AF, const SCEVUnknown *BP,

polly/trunk/lib/Analysis/ScopDetection.cpp

case llvm::Intrinsic::memset:
     return true;
   default:
     break;
   }
 
   return false;
 }
 
-bool ScopDetection::isInvariant(const Value &Val, const Region &Reg) const {
+bool ScopDetection::isInvariant(Value &Val, const Region &Reg,
+                                DetectionContext &Ctx) const {
   // A reference to function argument or constant value is invariant.
   if (isa<Argument>(Val) || isa<Constant>(Val))
     return true;
 
-  const Instruction *I = dyn_cast<Instruction>(&Val);
+  Instruction *I = dyn_cast<Instruction>(&Val);
   if (!I)
     return false;
 
   if (!Reg.contains(I))
     return true;
 
+  // Loads within the SCoP may read arbitrary values, need to hoist them. If it
+  // is not hoistable, it will be rejected later, but here we assume it is and
+  // that makes the value invariant.
+  if (auto LI = dyn_cast<LoadInst>(I)) {
+    Ctx.RequiredILS.insert(LI);
+    return true;
+  }
+
   if (I->mayHaveSideEffects())
     return false;
 
   if (isa<SelectInst>(I))
     return false;
 
   // When Val is a Phi node, it is likely not invariant. We do not check whether
   // Phi nodes are actually invariant, we assume that Phi nodes are usually not
   // invariant.
   if (isa<PHINode>(*I))
     return false;
 
   for (const Use &Operand : I->operands())
-    if (!isInvariant(*Operand, Reg))
+    if (!isInvariant(*Operand, Reg, Ctx))
       return false;
 
   return true;
 }
 
 /// Remove smax of smax(0, size) expressions from a SCEV expression and
 /// register the '...' components.
 ///
if (isa<UndefValue>(BV))
     return invalid<ReportUndefBasePtr>(Context, /*Assert=*/true, Inst);
 
   // FIXME: Think about allowing IntToPtrInst
   if (IntToPtrInst *Inst = dyn_cast<IntToPtrInst>(BV))
     return invalid<ReportIntToPtr>(Context, /*Assert=*/true, Inst);
 
   // Check that the base address of the access is invariant in the current
   // region.
-  if (!isInvariant(*BV, Context.CurRegion))
+  if (!isInvariant(*BV, Context.CurRegion, Context))
     return invalid<ReportVariantBasePtr>(Context, /*Assert=*/true, BV, Inst);
 
   AF = SE->getMinusSCEV(AF, BP);
 
   const SCEV *Size;
   if (!isa<MemIntrinsic>(Inst)) {
     Size = SE->getElementSize(Inst);
   } else {

polly/trunk/test/Isl/CodeGen/OpenMP/loop-body-references-outer-values-3.ll

-; RUN: opt %loadPolly -basicaa -polly-parallel -polly-parallel-force -polly-ast -analyze < %s | FileCheck %s -check-prefix=AST
-; RUN: opt %loadPolly -basicaa -polly-parallel -polly-parallel-force -polly-codegen -S -verify-dom-info < %s | FileCheck %s -check-prefix=IR
+; RUN: opt %loadPolly -basicaa -polly-parallel -polly-parallel-force -polly-invariant-load-hoisting=true -polly-ast -analyze < %s | FileCheck %s -check-prefix=AST
+; RUN: opt %loadPolly -basicaa -polly-parallel -polly-parallel-force -polly-invariant-load-hoisting=true -polly-codegen -S -verify-dom-info < %s | FileCheck %s -check-prefix=IR
 
 ; The interesting part of this test case is the instruction:
 ;   %tmp = bitcast i8* %call to i64**
 ; which is not part of the scop. In the SCEV based code generation not '%tmp',
 ; but %call is a parameter of the SCoP and we need to make sure its value is
 ; properly forwarded to the subfunction.
 
 ; AST: #pragma omp parallel for

polly/trunk/test/Isl/CodeGen/OpenMP/single_loop_with_loop_invariant_baseptr.ll

-; RUN: opt %loadPolly -tbaa -polly-parallel -polly-parallel-force -polly-parallel-force -polly-ast -analyze < %s | FileCheck %s -check-prefix=AST
-; RUN: opt %loadPolly -tbaa -polly-parallel -polly-parallel-force -polly-parallel-force -polly-codegen -S -verify-dom-info < %s | FileCheck %s -check-prefix=IR
+; RUN: opt %loadPolly -tbaa -polly-parallel -polly-parallel-force -polly-parallel-force -polly-invariant-load-hoisting=true -polly-ast -analyze < %s | FileCheck %s -check-prefix=AST
+; RUN: opt %loadPolly -tbaa -polly-parallel -polly-parallel-force -polly-parallel-force -polly-invariant-load-hoisting=true -polly-codegen -S -verify-dom-info < %s | FileCheck %s -check-prefix=IR
 
 ; #define N 1024
 ; float A[N];
 ;
 ; void single_parallel_loop(void) {
 ;   for (long i = 0; i < N; i++)
 ;     A[i] = 1;
 ; }

polly/trunk/test/ScopDetect/base_pointer.ll

-; RUN: opt %loadPolly -disable-basicaa -polly-detect -analyze < %s | FileCheck %s
+; RUN: opt %loadPolly -disable-basicaa -polly-detect -polly-invariant-load-hoisting=true -analyze < %s | FileCheck %s
 
 
 target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128"
 
 define void @base_pointer_in_condition(i64** noalias %A_ptr, i64 %N) nounwind {
 entry:
   fence seq_cst
   br label %pre

polly/trunk/test/ScopDetect/base_pointer_load_is_inst_inside_invariant_1___%for.i---%exit.jscop

+{
+   "arrays" : [
+      {
+         "name" : "MemRef_A",
+         "sizes" : [ "*" ],
+         "type" : "float*"
+      },
+      {
+         "name" : "MemRef_ptr",
+         "sizes" : [ "*" ],
+         "type" : "float"
+      }
+   ],
+   "context" : "[n] -> {  : -9223372036854775808 <= n <= 9223372036854775807 }",
+   "name" : "%for.i---%exit",
+   "statements" : [
+      {
+         "accesses" : [
+            {
+               "kind" : "read",
+               "relation" : "[n] -> { Stmt_S1[i0] -> MemRef_A[0] }"
+            },
+            {
+               "kind" : "write",
+               "relation" : "[n] -> { Stmt_S1[i0] -> MemRef_ptr[i0+1] }"
+            }
+         ],
+         "domain" : "[n] -> { Stmt_S1[i0] : 0 <= i0 < n }",
+         "name" : "Stmt_S1",
+         "schedule" : "[n] -> { Stmt_S1[i0] -> [i0] }"
+      }
+   ]
+}

polly/trunk/test/ScopDetect/base_pointer_load_setNewAccessRelation.ll

+; RUN: opt %loadPolly -polly-ignore-aliasing -polly-invariant-load-hoisting=true -polly-import-jscop -polly-import-jscop-dir=%S -polly-scops -polly-codegen -analyze < %s | FileCheck %s
+;
+; This violated an assertion in setNewAccessRelation that assumed base pointers
+; to be load-hoisted. Without this assertion, it codegen would generate invalid
+; code.
+;
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128"
+
+define void @base_pointer_load_is_inst_inside_invariant_1(i64 %n, float** %A) {
+entry:
+  br label %for.i
+
+for.i:
+  %indvar.i = phi i64 [ %indvar.i.next, %for.i.inc ], [ 0, %entry ]
+  br label %S1
+
+S1:
+  %ptr = load float*, float** %A
+  %conv = sitofp i64 %indvar.i to float
+  %arrayidx5 = getelementptr float, float* %ptr, i64 %indvar.i
+  store float %conv, float* %arrayidx5, align 4
+  br label %for.i.inc
+
+for.i.inc:
+  %indvar.i.next = add i64 %indvar.i, 1
+  %exitcond.i = icmp ne i64 %indvar.i.next, %n
+  br i1 %exitcond.i, label %for.i, label %exit
+
+exit:
+  ret void
+}
+
+
+; Detected by -polly-detect with required load hoist.
+; CHECK-NOT: Valid Region for Scop: for.i => exit
+;
+; Load hoist if %ptr by -polly-scops.
+; CHECK:      Invariant Accesses: {
+; CHECK-NEXT:     ReadAccess :=       [Reduction Type: NONE] [Scalar: 0]
+; CHECK-NEXT:         [n] -> { Stmt_S1[i0] -> MemRef_A[0] };
+; CHECK-NEXT:     Execution Context: [n] -> {  : n > 0 }
+; CHECK-NEXT: }

polly/trunk/test/ScopDetectionDiagnostics/ReportVariantBasePtr-01.ll

 ; RUN: opt %loadPolly -pass-remarks-missed="polly-detect" -polly-detect-track-failures -polly-detect -analyze < %s 2>&1| FileCheck %s
 
 ; struct b {
 ;   double **b;
 ; };
 ;
 ; void a(struct b *A) {
 ;   for (int i=0; i<32; i++)
 ;     A[i].b[i] = 0;
 ; }
 
+; The loads are currently just adds %7 to the list of required invariant loads
+; and only -polly-scops checks whether it is actionally possible the be load
+; hoisted. The SCoP is still rejected by -polly-detect because it may alias
+; with %A and is not considered to be eligble for runtime alias checking.
+
 ; CHECK: remark: ReportVariantBasePtr01.c:6:8: The following errors keep this region from being a Scop.
-; CHECK: remark: ReportVariantBasePtr01.c:7:5: The base address of this array is not invariant inside the loop
+; CHECK: remark: ReportVariantBasePtr01.c:7:5: Accesses to the arrays "A", " <unknown> " may access the same memory.
 
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 %struct.b = type { double** }
 
 define void @a(%struct.b* nocapture readonly %A) #0 !dbg !4 {
 entry:
   br label %entry.split

polly/trunk/test/ScopInfo/ranged_parameter_2.ll

-; RUN: opt %loadPolly -polly-scops -analyze -polly-allow-nonaffine < %s \
+; RUN: opt %loadPolly -polly-scops -analyze -polly-allow-nonaffine -polly-invariant-load-hoisting=true < %s \
 ; RUN:  -debug 2>&1 | FileCheck %s
 
 ; REQUIRES: asserts
 
 ; CHECK: Region: %bb1---%bb16
 ; CHECK:   [n] -> {  : 1 = 0 }
 
 ; This test case at some point caused an assertion when modeling a scop, due

polly/trunk/test/ScopInfo/variant_base_pointer.ll

-; RUN: opt %loadPolly -polly-ignore-aliasing -polly-scops -analyze < %s | FileCheck %s
-; RUN: opt %loadPolly -polly-ignore-aliasing -polly-codegen -analyze < %s
+; RUN: opt %loadPolly -polly-ignore-aliasing -polly-invariant-load-hoisting=true -polly-scops -analyze < %s | FileCheck %s
+; RUN: opt %loadPolly -polly-ignore-aliasing -polly-invariant-load-hoisting=true -polly-codegen -analyze < %s
 ;
-; CHECK:       Invariant Accesses: {
-; CHECK-NEXT:  }
+; %tmp is added to the list of required hoists by -polly-scops and just
+; assumed to be hoisted. Only -polly-scops recognizes it to be unhoistable
+; because ir depends on %call which cannot be executed speculatively.
+;
+; CHECK-NOT:       Invariant Accesses:
 ;
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 ; Function Attrs: nounwind uwtable
 define void @cli_hex2int() {
 entry:
   br label %if.end