[Polly] Compute and print the loop carried dependency distance

[jdoerfert]

Summary

During the IslAst parallelism check also compute the actual dependency
distance and store it in the IstAst for node.

[jdoerfert]

ping

[jdoerfert]

Rebase patch and show unit tests result in phabricator

[sebpop]

lib/Analysis/Dependences.cpp

521

What happens when the isl_set_free function returns a non-zero error?

test/Isl/Ast/dependence_distance_parametric.ll

5

j++)

test/Isl/Ast/dependence_distance_varying.ll

4

I am not sure I understand what "((N - 1) % 2) + 1" means: for a given N, say 100, this expression is equal to 2,

for (int i = 0; i < N; i++)
  A[i] = A[N - i] + 1;

and the dependence distance in this loop is not constant:

• A[1] is written at iteration i=1 and is read at iteration i=99 -> distance 98
• A[2] is written at iteration i=2 and is read at iteration i=98 -> distance 96
• etc.

test/Isl/Ast/dependence_distance_varying_in_outer_loop.ll

6

Is the distance carried by the inner loop equal to 1?

Say we fix the outer loop j to iteration 5, then in the inner loop

• first iteration i=5, A[2] is written and A[5] is read,
• third iteration i=8, A[5] is written and A[5] is read.

I think that the dependence distance is not constant in this case.

8

for j < 3, A[i - 3] is a negative access function.

test/Isl/Ast/dependence_distance_varying_multiple.ll

5

The difference between A[i] and A[100 - 2*i] is not constant.

This loop does not carry a constant dependence distance of 1, 2, or 5.

[jdoerfert]

Answers

lib/Analysis/Dependences.cpp

521

isl "guarantees" (with the __isl_null annotation) that we get a nullptr back.

Apperently ppl don't like my shortcut, I can change that.

test/Isl/Ast/dependence_distance_parametric.ll

5

thx

test/Isl/Ast/dependence_distance_varying.ll

4

We compute the minimal dependence distance, everything else seems not really useful? Am I wrong here?

test/Isl/Ast/dependence_distance_varying_in_outer_loop.ll

6

Again, 1 is the minimal dependence distance

8

That is totally fine for a subset of possible pointers A.
What is the problem with negative access functions?

test/Isl/Ast/dependence_distance_varying_multiple.ll

5

Should again be the minimal dep. distance.

[sebpop]

test/Isl/Ast/dependence_distance_varying_multiple.ll

5

Ok, so then let's print:

#pragma minimal dependence distance

that is different than what books call "dependence distance", that's why I got confused.

[jdoerfert]

test/Isl/Ast/dependence_distance_varying_multiple.ll

5

True, will do.

[jdoerfert]

Addressed review concerns

[jdoerfert]

@sebpop Is it good to go like this?

[sebpop]

LGTM with those changes.
Let's also wait for Tobi's ok before commit.

include/polly/CodeGen/IslAst.h

55

MinimalDependenceDistance

72

the minimal dependence distance

73

MinimalDependenceDistance

132

Get the minimal dependence distance or nullptr.

134

getMinimalDependenceDistance

include/polly/Dependences.h

95

MinDistancePtr If not nullptr, the minimal dependence distance will be returned at the address of that pointer

lib/Analysis/Dependences.cpp

490

MinDistancePtr

527

Please add a comment above this stmt to explain that we take the min of the distance polyhedron.

[jdoerfert]

Closed by commit rL217729 (authored by @jdoerfert).

include/polly/CodeGen/IslAst.h

public:
   using MemoryAccessSet = SmallPtrSet<MemoryAccess *, 4>;
 
   /// @brief Payload information used to annotate an AST node.
   struct IslAstUserPayload {
     /// @brief Construct and initialize the payload.
     IslAstUserPayload()
         : IsInnermost(false), IsInnermostParallel(false),
           IsOutermostParallel(false), IsReductionParallel(false),
-          Build(nullptr) {}
+          DependenceDistance(nullptr), Build(nullptr) {}
 
     /// @brief Cleanup all isl structs on destruction.
     ~IslAstUserPayload();
 
     /// @brief Flag to mark innermost loops.
     bool IsInnermost;
 
     /// @brief Flag to mark innermost parallel loops.
     bool IsInnermostParallel;
 
     /// @brief Flag to mark outermost parallel loops.
     bool IsOutermostParallel;
 
     /// @brief Flag to mark parallel loops which break reductions.
     bool IsReductionParallel;
 
+    /// @brief The dependence distance for this loop if constant and not zero.
+    struct isl_pw_aff *DependenceDistance;
+
     /// @brief The build environment at the time this node was constructed.
     isl_ast_build *Build;
 
     /// @brief Set of accesses which break reduction dependences.
     MemoryAccessSet BrokenReductions;
   };
 
 private:
public:
   static bool isInnermostParallel(__isl_keep isl_ast_node *Node);
 
   /// @brief Is this loop a reduction parallel loop?
   static bool isReductionParallel(__isl_keep isl_ast_node *Node);
 
   /// @brief Get the nodes schedule or a nullptr if not available.
   static __isl_give isl_union_map *getSchedule(__isl_keep isl_ast_node *Node);
 
+  /// @brief Get the nodes dependence distance or a nullptr if not available.
+  static __isl_give isl_pw_aff *
+  getDependenceDistance(__isl_keep isl_ast_node *Node);
+
   /// @brief Get the nodes broken reductions or a nullptr if not available.
   static MemoryAccessSet *getBrokenReductions(__isl_keep isl_ast_node *Node);
 
   /// @brief Get the nodes build context or a nullptr if not available.
   static __isl_give isl_ast_build *getBuild(__isl_keep isl_ast_node *Node);
 
   ///}
 

include/polly/Dependences.h

 #ifndef POLLY_DEPENDENCES_H
 #define POLLY_DEPENDENCES_H
 
 #include "polly/ScopPass.h"
 
 #include <map>
 #include "isl/ctx.h"
 
+struct isl_pw_aff;
 struct isl_union_map;
 struct isl_union_set;
 struct isl_map;
 struct isl_set;
 struct clast_for;
 
 using namespace llvm;
 
public:
   /// @param NewScattering The new scatterings
   ///
   /// @return bool True if the new scattering is valid, false it it reverses
   ///              dependences.
   bool isValidScattering(StatementToIslMapTy *NewScatterings);
 
   /// @brief Check if a partial schedule is parallel wrt to @p Deps.
   ///
-  /// @param Schedule The subset of the scattering space that we want to check.
-  /// @param Deps     The dependences @p Schedule needs to respect.
+  /// @param Schedule    The subset of the scattering space that we want to
+  ///                    check.
+  /// @param Deps        The dependences @p Schedule needs to respect.
+  /// @param DistancePtr If given the dependences distance will be return here.
   ///
   /// @return Returns true, if executing parallel the outermost dimension of
   ///         @p Schedule is valid according to the dependences @p Deps.
   bool isParallel(__isl_keep isl_union_map *Schedule,
-                  __isl_take isl_union_map *Deps);
+                  __isl_take isl_union_map *Deps,
+                  __isl_give isl_pw_aff **DistancePtr = nullptr);
 
   /// @brief Get the dependences in this Scop.
   ///
   /// @param Kinds This integer defines the different kinds of dependences
   ///              that will be returned. To return more than one kind, the
   ///              different kinds are 'ored' together.
   isl_union_map *getDependences(int Kinds);
 

lib/Analysis/Dependences.cpp

 // Parallelism test: if the distance is zero in all outer dimensions, then it
 // has to be zero in the current dimension as well.
 //
 // Implementation: first, translate dependences into time space, then force
 // outer dimensions to be equal. If the distance is zero in the current
 // dimension, then the loop is parallel. The distance is zero in the current
 // dimension if it is a subset of a map with equal values for the current
 // dimension.
-bool Dependences::isParallel(isl_union_map *Schedule, isl_union_map *Deps) {
-  isl_map *ScheduleDeps, *Test;
-  unsigned Dimension, IsParallel;
+bool Dependences::isParallel(isl_union_map *Schedule, isl_union_map *Deps,
+                             isl_pw_aff **DistancePtr) {
+  isl_set *Deltas, *Distance;
+  isl_map *ScheduleDeps;
+  unsigned Dimension;
+  bool IsParallel;
 
   Deps = isl_union_map_apply_range(Deps, isl_union_map_copy(Schedule));
   Deps = isl_union_map_apply_domain(Deps, isl_union_map_copy(Schedule));
 
   if (isl_union_map_is_empty(Deps)) {
     isl_union_map_free(Deps);
     return true;
   }
 
   ScheduleDeps = isl_map_from_union_map(Deps);
   Dimension = isl_map_dim(ScheduleDeps, isl_dim_out) - 1;
 
   for (unsigned i = 0; i < Dimension; i++)
     ScheduleDeps = isl_map_equate(ScheduleDeps, isl_dim_out, i, isl_dim_in, i);
 
-  Test = isl_map_universe(isl_map_get_space(ScheduleDeps));
-  Test = isl_map_equate(Test, isl_dim_out, Dimension, isl_dim_in, Dimension);
-  IsParallel = isl_map_is_subset(ScheduleDeps, Test);
+  Deltas = isl_map_deltas(ScheduleDeps);
+  Distance = isl_set_universe(isl_set_get_space(Deltas));
 
-  isl_map_free(Test);
-  isl_map_free(ScheduleDeps);
+  // [0, ..., 0, +] - All zeros and last dimension larger than zero
+  for (unsigned i = 0; i < Dimension; i++)
+    Distance = isl_set_fix_si(Distance, isl_dim_set, i, 0);
+
+  Distance = isl_set_lower_bound_si(Distance, isl_dim_set, Dimension, 1);
+  Distance = isl_set_intersect(Distance, Deltas);
 
-  return IsParallel;
+  IsParallel = isl_set_is_empty(Distance);
+  if (IsParallel && !isl_set_free(Distance))
+    return true;
+
+  if (!DistancePtr && !isl_set_free(Distance))
+    return false;
+
+  Distance = isl_set_project_out(Distance, isl_dim_set, 0, Dimension);
+  Distance = isl_set_coalesce(Distance);
+
+  *DistancePtr = isl_pw_aff_coalesce(isl_set_dim_min(Distance, 0));
+
+  return false;
 }
 
 static void printDependencyMap(raw_ostream &OS, __isl_keep isl_union_map *DM) {
   if (DM)
     OS << DM << "\n";
   else
     OS << "n/a\n";
 }

lib/CodeGen/IslAst.cpp

 
 /// @brief Free an IslAstUserPayload object pointed to by @p Ptr
 static void freeIslAstUserPayload(void *Ptr) {
   delete ((IslAstInfo::IslAstUserPayload *)Ptr);
 }
 
 IslAstInfo::IslAstUserPayload::~IslAstUserPayload() {
   isl_ast_build_free(Build);
+  isl_pw_aff_free(DependenceDistance);
 }
 
 /// @brief Temporary information used when building the ast.
 struct AstBuildUserInfo {
   /// @brief Construct and initialize the helper struct for AST creation.
   AstBuildUserInfo()
       : Deps(nullptr), InParallelFor(false), LastForNodeId(nullptr) {}
 
   /// @brief The dependence information used for the parallelism check.
   Dependences *Deps;
 
   /// @brief Flag to indicate that we are inside a parallel for node.
   bool InParallelFor;
 
   /// @brief The last iterator id created for the current SCoP.
   isl_id *LastForNodeId;
 };
 
 /// @brief Print a string @p str in a single line using @p Printer.
 static isl_printer *printLine(__isl_take isl_printer *Printer,
-                              const std::string &str) {
+                              const std::string &str,
+                              __isl_keep isl_pw_aff *PWA = nullptr) {
   Printer = isl_printer_start_line(Printer);
   Printer = isl_printer_print_str(Printer, str.c_str());
+  if (PWA)
+    Printer = isl_printer_print_pw_aff(Printer, PWA);
   return isl_printer_end_line(Printer);
 }
 
 /// @brief Return all broken reductions as a string of clauses (OpenMP style).
 static const std::string getBrokenReductionsStr(__isl_keep isl_ast_node *Node) {
   IslAstInfo::MemoryAccessSet *BrokenReductions;
   std::string str;
 
static const std::string getBrokenReductionsStr(__isl_keep isl_ast_node *Node) {
   return str;
 }
 
 /// @brief Callback executed for each for node in the ast in order to print it.
 static isl_printer *cbPrintFor(__isl_take isl_printer *Printer,
                                __isl_take isl_ast_print_options *Options,
                                __isl_keep isl_ast_node *Node, void *) {
 
+  isl_pw_aff *DD = IslAstInfo::getDependenceDistance(Node);
   const std::string BrokenReductionsStr = getBrokenReductionsStr(Node);
+  const std::string DepDisPragmaStr = "#pragma dependence distance: ";
   const std::string SimdPragmaStr = "#pragma simd";
   const std::string OmpPragmaStr = "#pragma omp parallel for";
 
+  if (DD)
+    Printer = printLine(Printer, DepDisPragmaStr, DD);
+
   if (IslAstInfo::isInnermostParallel(Node))
     Printer = printLine(Printer, SimdPragmaStr + BrokenReductionsStr);
 
   if (IslAstInfo::isOutermostParallel(Node))
     Printer = printLine(Printer, OmpPragmaStr + BrokenReductionsStr);
 
+  isl_pw_aff_free(DD);
   return isl_ast_node_for_print(Node, Printer, Options);
 }
 
 /// @brief Check if the current scheduling dimension is parallel
 ///
 /// In case the dimension is parallel we also check if any reduction
 /// dependences is broken when we exploit this parallelism. If so,
 /// @p IsReductionParallel will be set to true. The reduction dependences we use
 /// to check are actually the union of the transitive closure of the initial
 /// reduction dependences together with their reveresal. Even though these
 /// dependences connect all iterations with each other (thus they are cyclic)
 /// we can perform the parallelism check as we are only interested in a zero
 /// (or non-zero) dependence distance on the dimension in question.
 static bool astScheduleDimIsParallel(__isl_keep isl_ast_build *Build,
                                      Dependences *D,
                                      IslAstUserPayload *NodeInfo) {
   if (!D->hasValidDependences())
     return false;
 
   isl_union_map *Schedule = isl_ast_build_get_schedule(Build);
   isl_union_map *Deps = D->getDependences(
       Dependences::TYPE_RAW | Dependences::TYPE_WAW | Dependences::TYPE_WAR);
-  if (!D->isParallel(Schedule, Deps) && !isl_union_map_free(Schedule))
+
+  if (!D->isParallel(Schedule, Deps, &NodeInfo->DependenceDistance) &&
+      !isl_union_map_free(Schedule))
     return false;
 
   isl_union_map *RedDeps = D->getDependences(Dependences::TYPE_TC_RED);
   if (!D->isParallel(Schedule, RedDeps))
     NodeInfo->IsReductionParallel = true;
 
   if (!NodeInfo->IsReductionParallel && !isl_union_map_free(Schedule))
     return true;
bool IslAstInfo::isReductionParallel(__isl_keep isl_ast_node *Node) {
   return Payload && Payload->IsReductionParallel;
 }
 
 isl_union_map *IslAstInfo::getSchedule(__isl_keep isl_ast_node *Node) {
   IslAstUserPayload *Payload = getNodePayload(Node);
   return Payload ? isl_ast_build_get_schedule(Payload->Build) : nullptr;
 }
 
+isl_pw_aff *IslAstInfo::getDependenceDistance(__isl_keep isl_ast_node *Node) {
+  IslAstUserPayload *Payload = getNodePayload(Node);
+  return Payload ? isl_pw_aff_copy(Payload->DependenceDistance) : nullptr;
+}
+
 IslAstInfo::MemoryAccessSet *
 IslAstInfo::getBrokenReductions(__isl_keep isl_ast_node *Node) {
   IslAstUserPayload *Payload = getNodePayload(Node);
   return Payload ? &Payload->BrokenReductions : nullptr;
 }
 
 isl_ast_build *IslAstInfo::getBuild(__isl_keep isl_ast_node *Node) {
   IslAstUserPayload *Payload = getNodePayload(Node);

test/Isl/Ast/dependence_distance_constant.ll

+; RUN: opt %loadPolly -polly-ast -polly-ast-detect-parallel -analyze < %s | FileCheck %s
+;
+;        void f(int *A, int N) {
+; CHECK:   #pragma dependence distance: 1
+;          for (int j = 0; j < N; j++)
+; CHECK:     #pragma dependence distance: 8
+;            for (int i = 0; i < N; i++)
+;              A[i + 8] = A[i] + 1;
+;        }
+;
+target datalayout = "e-m:e-p:32:32-i64:64-v128:64:128-n32-S64"
+
+define void @f(i32* %A, i32 %N) {
+entry:
+  br label %for.cond
+
+for.cond:                                         ; preds = %for.inc6, %entry
+  %j.0 = phi i32 [ 0, %entry ], [ %inc7, %for.inc6 ]
+  %cmp = icmp slt i32 %j.0, %N
+  br i1 %cmp, label %for.body, label %for.end8
+
+for.body:                                         ; preds = %for.cond
+  br label %for.cond1
+
+for.cond1:                                        ; preds = %for.inc, %for.body
+  %i.0 = phi i32 [ 0, %for.body ], [ %inc, %for.inc ]
+  %exitcond = icmp ne i32 %i.0, %N
+  br i1 %exitcond, label %for.body3, label %for.end
+
+for.body3:                                        ; preds = %for.cond1
+  %arrayidx = getelementptr inbounds i32* %A, i32 %i.0
+  %tmp = load i32* %arrayidx, align 4
+  %add = add nsw i32 %tmp, 1
+  %add4 = add nsw i32 %i.0, 8
+  %arrayidx5 = getelementptr inbounds i32* %A, i32 %add4
+  store i32 %add, i32* %arrayidx5, align 4
+  br label %for.inc
+
+for.inc:                                          ; preds = %for.body3
+  %inc = add nsw i32 %i.0, 1
+  br label %for.cond1
+
+for.end:                                          ; preds = %for.cond1
+  br label %for.inc6
+
+for.inc6:                                         ; preds = %for.end
+  %inc7 = add nsw i32 %j.0, 1
+  br label %for.cond
+
+for.end8:                                         ; preds = %for.cond
+  ret void
+}

test/Isl/Ast/dependence_distance_multiple_constant.ll

+; RUN: opt %loadPolly -basicaa -polly-ast -polly-ast-detect-parallel -analyze < %s | FileCheck %s
+;
+;        void f(int *restrict A, int *restrict B, int N) {
+; CHECK:   #pragma dependence distance: 5
+;          for (int i = 0; i < N; i++) {
+;            A[i + 7] = A[i] + 1;
+;            B[i + 5] = B[i] + 1;
+;          }
+;        }
+;
+target datalayout = "e-m:e-p:32:32-i64:64-v128:64:128-n32-S64"
+
+define void @f(i32* noalias %A, i32* noalias %B, i32 %N) {
+entry:
+  br label %for.cond
+
+for.cond:                                         ; preds = %for.inc, %entry
+  %i.0 = phi i32 [ 0, %entry ], [ %inc, %for.inc ]
+  %cmp = icmp slt i32 %i.0, %N
+  br i1 %cmp, label %for.body, label %for.end
+
+for.body:                                         ; preds = %for.cond
+  %arrayidx = getelementptr inbounds i32* %A, i32 %i.0
+  %tmp = load i32* %arrayidx, align 4
+  %add = add nsw i32 %tmp, 1
+  %add1 = add nsw i32 %i.0, 7
+  %arrayidx2 = getelementptr inbounds i32* %A, i32 %add1
+  store i32 %add, i32* %arrayidx2, align 4
+  %arrayidx3 = getelementptr inbounds i32* %B, i32 %i.0
+  %tmp1 = load i32* %arrayidx3, align 4
+  %add4 = add nsw i32 %tmp1, 1
+  %add5 = add nsw i32 %i.0, 5
+  %arrayidx6 = getelementptr inbounds i32* %B, i32 %add5
+  store i32 %add4, i32* %arrayidx6, align 4
+  br label %for.inc
+
+for.inc:                                          ; preds = %for.body
+  %inc = add nsw i32 %i.0, 1
+  br label %for.cond
+
+for.end:                                          ; preds = %for.cond
+  ret void
+}

test/Isl/Ast/dependence_distance_parametric.ll

+; RUN: opt %loadPolly -polly-ast -polly-ast-detect-parallel -analyze < %s | FileCheck %s
+;
+;        void f(int *A, int N, int c) {
+; CHECK:   #pragma dependence distance: 1
+;          for (int j = 0; j < N; j+++
+; CHECK:     #pragma dependence distance: (-1 + c >= 0) ? (c) : -c
+;            for (int i = 0; i < N; i++)
+;              A[i + c] = A[i] + 1;
+;        }
+;
+target datalayout = "e-m:e-p:32:32-i64:64-v128:64:128-n32-S64"
+
+define void @f(i32* %A, i32 %N, i32 %c) {
+entry:
+  br label %for.cond
+
+for.cond:                                         ; preds = %for.inc6, %entry
+  %j.0 = phi i32 [ 0, %entry ], [ %inc7, %for.inc6 ]
+  %cmp = icmp slt i32 %j.0, %N
+  br i1 %cmp, label %for.body, label %for.end8
+
+for.body:                                         ; preds = %for.cond
+  br label %for.cond1
+
+for.cond1:                                        ; preds = %for.inc, %for.body
+  %i.0 = phi i32 [ 0, %for.body ], [ %inc, %for.inc ]
+  %exitcond = icmp ne i32 %i.0, %N
+  br i1 %exitcond, label %for.body3, label %for.end
+
+for.body3:                                        ; preds = %for.cond1
+  %arrayidx = getelementptr inbounds i32* %A, i32 %i.0
+  %tmp = load i32* %arrayidx, align 4
+  %add = add nsw i32 %tmp, 1
+  %add4 = add nsw i32 %i.0, %c
+  %arrayidx5 = getelementptr inbounds i32* %A, i32 %add4
+  store i32 %add, i32* %arrayidx5, align 4
+  br label %for.inc
+
+for.inc:                                          ; preds = %for.body3
+  %inc = add nsw i32 %i.0, 1
+  br label %for.cond1
+
+for.end:                                          ; preds = %for.cond1
+  br label %for.inc6
+
+for.inc6:                                         ; preds = %for.end
+  %inc7 = add nsw i32 %j.0, 1
+  br label %for.cond
+
+for.end8:                                         ; preds = %for.cond
+  ret void
+}

test/Isl/Ast/dependence_distance_parametric_expr.ll

+; RUN: opt %loadPolly -polly-ast -polly-ast-detect-parallel -analyze < %s | FileCheck %s
+;
+;        void f(int *A, int N, int c, int v) {
+; CHECK:   #pragma dependence distance: 1
+;          for (int j = 0; j < N; j++)
+; CHECK:     #pragma dependence distance: (-1 + c + v >= 0) ? (c + v) : -c - v
+;            for (int i = 0; i < N; i++)
+;              A[i + c + v] = A[i] + 1;
+;        }
+;
+target datalayout = "e-m:e-p:32:32-i64:64-v128:64:128-n32-S64"
+
+define void @f(i32* %A, i32 %N, i32 %c, i32 %v) {
+entry:
+  br label %for.cond
+
+for.cond:                                         ; preds = %for.inc7, %entry
+  %j.0 = phi i32 [ 0, %entry ], [ %inc8, %for.inc7 ]
+  %cmp = icmp slt i32 %j.0, %N
+  br i1 %cmp, label %for.body, label %for.end9
+
+for.body:                                         ; preds = %for.cond
+  br label %for.cond1
+
+for.cond1:                                        ; preds = %for.inc, %for.body
+  %i.0 = phi i32 [ 0, %for.body ], [ %inc, %for.inc ]
+  %exitcond = icmp ne i32 %i.0, %N
+  br i1 %exitcond, label %for.body3, label %for.end
+
+for.body3:                                        ; preds = %for.cond1
+  %arrayidx = getelementptr inbounds i32* %A, i32 %i.0
+  %tmp = load i32* %arrayidx, align 4
+  %add = add nsw i32 %tmp, 1
+  %add4 = add nsw i32 %i.0, %c
+  %add5 = add nsw i32 %add4, %v
+  %arrayidx6 = getelementptr inbounds i32* %A, i32 %add5
+  store i32 %add, i32* %arrayidx6, align 4
+  br label %for.inc
+
+for.inc:                                          ; preds = %for.body3
+  %inc = add nsw i32 %i.0, 1
+  br label %for.cond1
+
+for.end:                                          ; preds = %for.cond1
+  br label %for.inc7
+
+for.inc7:                                         ; preds = %for.end
+  %inc8 = add nsw i32 %j.0, 1
+  br label %for.cond
+
+for.end9:                                         ; preds = %for.cond
+  ret void
+}

test/Isl/Ast/dependence_distance_varying.ll

+; RUN: opt %loadPolly -polly-ast -polly-ast-detect-parallel -analyze < %s | FileCheck %s
+;
+;         void f(int *A, int N) {
+; CHECK:    #pragma dependence distance: ((N - 1) % 2) + 1
+;           for (int i = 0; i < N; i++)
+;             A[i] = A[N - i] + 1;
+;         }
+;
+target datalayout = "e-m:e-p:32:32-i64:64-v128:64:128-n32-S64"
+
+define void @f(i32* %A, i32 %N) {
+entry:
+  br label %for.cond
+
+for.cond:                                         ; preds = %for.inc, %entry
+  %i.0 = phi i32 [ 0, %entry ], [ %inc, %for.inc ]
+  %cmp = icmp slt i32 %i.0, %N
+  br i1 %cmp, label %for.body, label %for.end
+
+for.body:                                         ; preds = %for.cond
+  %sub = sub nsw i32 %N, %i.0
+  %arrayidx = getelementptr inbounds i32* %A, i32 %sub
+  %tmp = load i32* %arrayidx, align 4
+  %add = add nsw i32 %tmp, 1
+  %arrayidx1 = getelementptr inbounds i32* %A, i32 %i.0
+  store i32 %add, i32* %arrayidx1, align 4
+  br label %for.inc
+
+for.inc:                                          ; preds = %for.body
+  %inc = add nsw i32 %i.0, 1
+  br label %for.cond
+
+for.end:                                          ; preds = %for.cond
+  ret void
+}

test/Isl/Ast/dependence_distance_varying_in_outer_loop.ll

+; RUN: opt %loadPolly -polly-canonicalize -polly-ast -polly-ast-detect-parallel -analyze < %s | FileCheck %s
+;
+;        void f(int *restrict A, int *restrict sum) {
+; CHECK: #pragma dependence distance: 1
+;          for (int j = 0; j < 1024; j++)
+; CHECK: #pragma dependence distance: 1
+;            for (int i = j; i < 1024; i++)
+;              A[i - 3] = A[j] * 2 + A[j] + 2;
+;        }
+;
+target datalayout = "e-m:e-p:32:32-i64:64-v128:64:128-n32-S64"
+
+define void @f(i32* noalias %A, i32* noalias %sum) {
+entry:
+  br label %for.cond
+
+for.cond:                                         ; preds = %for.inc7, %entry
+  %j.0 = phi i32 [ 0, %entry ], [ %inc8, %for.inc7 ]
+  %exitcond1 = icmp ne i32 %j.0, 1024
+  br i1 %exitcond1, label %for.body, label %for.end9
+
+for.body:                                         ; preds = %for.cond
+  br label %for.cond1
+
+for.cond1:                                        ; preds = %for.inc, %for.body
+  %i.0 = phi i32 [ %j.0, %for.body ], [ %inc, %for.inc ]
+  %exitcond = icmp ne i32 %i.0, 1024
+  br i1 %exitcond, label %for.body3, label %for.end
+
+for.body3:                                        ; preds = %for.cond1
+  %arrayidx = getelementptr inbounds i32* %A, i32 %j.0
+  %tmp = load i32* %arrayidx, align 4
+  %add = mul nsw i32 %tmp, 3
+  %add5 = add nsw i32 %add, 2
+  %sub = add nsw i32 %i.0, -3
+  %arrayidx6 = getelementptr inbounds i32* %A, i32 %sub
+  store i32 %add5, i32* %arrayidx6, align 4
+  br label %for.inc
+
+for.inc:                                          ; preds = %for.body3
+  %inc = add nsw i32 %i.0, 1
+  br label %for.cond1
+
+for.end:                                          ; preds = %for.cond1
+  br label %for.inc7
+
+for.inc7:                                         ; preds = %for.end
+  %inc8 = add nsw i32 %j.0, 1
+  br label %for.cond
+
+for.end9:                                         ; preds = %for.cond
+  ret void
+}

test/Isl/Ast/dependence_distance_varying_multiple.ll

+; RUN: opt %loadPolly -basicaa -polly-ast -polly-ast-detect-parallel -analyze < %s | FileCheck %s
+;
+;        void f(int *restrict A, int *restrict B, int *restrict C, int *restrict D,
+;               int *restrict E, int N) {
+; CHECK:   #pragma dependence distance: (-35 + N >= 0) ? (1) : (-17 + N >= 0 && 34 - N >= 0) ? (2) : 5
+;          for (int i = 0; i < N; i++) {
+;            A[i] = A[100 - 2 * i] + 1;
+;            B[i] = B[100 - 3 * i] + 1;
+;            C[i] = C[100 - 4 * i] + 1;
+;            D[i] = D[100 - 5 * i] + 1;
+;            E[i] = E[100 - 6 * i] + 1;
+;          }
+;        }
+;
+target datalayout = "e-m:e-p:32:32-i64:64-v128:64:128-n32-S64"
+
+define void @f(i32* noalias %A, i32* noalias %B, i32* noalias %C, i32* noalias %D, i32* noalias %E, i32 %N) {
+entry:
+  br label %for.cond
+
+for.cond:                                         ; preds = %for.inc, %entry
+  %i.0 = phi i32 [ 0, %entry ], [ %inc, %for.inc ]
+  %cmp = icmp slt i32 %i.0, %N
+  br i1 %cmp, label %for.body, label %for.end
+
+for.body:                                         ; preds = %for.cond
+  %mul = shl nsw i32 %i.0, 1
+  %sub = sub nsw i32 100, %mul
+  %arrayidx = getelementptr inbounds i32* %A, i32 %sub
+  %tmp = load i32* %arrayidx, align 4
+  %add = add nsw i32 %tmp, 1
+  %arrayidx1 = getelementptr inbounds i32* %A, i32 %i.0
+  store i32 %add, i32* %arrayidx1, align 4
+  %tmp1 = mul i32 %i.0, -3
+  %sub3 = add i32 %tmp1, 100
+  %arrayidx4 = getelementptr inbounds i32* %B, i32 %sub3
+  %tmp2 = load i32* %arrayidx4, align 4
+  %add5 = add nsw i32 %tmp2, 1
+  %arrayidx6 = getelementptr inbounds i32* %B, i32 %i.0
+  store i32 %add5, i32* %arrayidx6, align 4
+  %mul7 = shl nsw i32 %i.0, 2
+  %sub8 = sub nsw i32 100, %mul7
+  %arrayidx9 = getelementptr inbounds i32* %C, i32 %sub8
+  %tmp3 = load i32* %arrayidx9, align 4
+  %add10 = add nsw i32 %tmp3, 1
+  %arrayidx11 = getelementptr inbounds i32* %C, i32 %i.0
+  store i32 %add10, i32* %arrayidx11, align 4
+  %tmp4 = mul i32 %i.0, -5
+  %sub13 = add i32 %tmp4, 100
+  %arrayidx14 = getelementptr inbounds i32* %D, i32 %sub13
+  %tmp5 = load i32* %arrayidx14, align 4
+  %add15 = add nsw i32 %tmp5, 1
+  %arrayidx16 = getelementptr inbounds i32* %D, i32 %i.0
+  store i32 %add15, i32* %arrayidx16, align 4
+  %tmp6 = mul i32 %i.0, -6
+  %sub18 = add i32 %tmp6, 100
+  %arrayidx19 = getelementptr inbounds i32* %E, i32 %sub18
+  %tmp7 = load i32* %arrayidx19, align 4
+  %add20 = add nsw i32 %tmp7, 1
+  %arrayidx21 = getelementptr inbounds i32* %E, i32 %i.0
+  store i32 %add20, i32* %arrayidx21, align 4
+  br label %for.inc
+
+for.inc:                                          ; preds = %for.body
+  %inc = add nsw i32 %i.0, 1
+  br label %for.cond
+
+for.end:                                          ; preds = %for.cond
+  ret void
+}

test/Isl/Ast/reduction_dependences_equal_non_reduction_dependences.ll

 ; RUN: opt %loadPolly -basicaa -polly-ast -polly-ast-detect-parallel -analyze < %s | FileCheck %s
 ;
 ; This loopnest contains a reduction which imposes the same dependences as the
 ; accesses to the array A. We need to ensure we do __not__ parallelize anything
 ; here.
 ;
+; CHECK: pragma dependence distance: 1
 ; CHECK-NOT: pragma
 ; CHECK-NOT: reduction
 ;
 ;    void AandSum(int *restrict sum, int *restrict A) {
 ;      for (int i = 0; i < 1024; i++) {
 ;        A[i] = A[i] + A[i - 1];
 ;        A[i - 1] = A[i] + A[i - 2];
 ;        *sum += i;