2012-10-03 03:50:38 +08:00
|
|
|
//===- IslAst.cpp - isl code generator interface --------------------------===//
|
|
|
|
//
|
|
|
|
// The LLVM Compiler Infrastructure
|
|
|
|
//
|
|
|
|
// This file is distributed under the University of Illinois Open Source
|
|
|
|
// License. See LICENSE.TXT for details.
|
|
|
|
//
|
|
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
//
|
|
|
|
// The isl code generator interface takes a Scop and generates a isl_ast. This
|
|
|
|
// ist_ast can either be returned directly or it can be pretty printed to
|
|
|
|
// stdout.
|
|
|
|
//
|
|
|
|
// A typical isl_ast output looks like this:
|
|
|
|
//
|
|
|
|
// for (c2 = max(0, ceild(n + m, 2); c2 <= min(511, floord(5 * n, 3)); c2++) {
|
|
|
|
// bb2(c2);
|
|
|
|
// }
|
|
|
|
//
|
|
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
|
2012-12-18 15:46:13 +08:00
|
|
|
#include "polly/CodeGen/CodeGeneration.h"
|
2012-10-03 03:50:38 +08:00
|
|
|
#include "polly/CodeGen/IslAst.h"
|
2012-12-13 14:24:06 +08:00
|
|
|
#include "polly/Dependences.h"
|
2013-05-07 16:11:54 +08:00
|
|
|
#include "polly/LinkAllPasses.h"
|
2013-05-07 15:31:10 +08:00
|
|
|
#include "polly/Options.h"
|
2012-10-03 03:50:38 +08:00
|
|
|
#include "polly/ScopInfo.h"
|
|
|
|
#include "llvm/Support/Debug.h"
|
|
|
|
|
|
|
|
#include "isl/union_map.h"
|
|
|
|
#include "isl/list.h"
|
|
|
|
#include "isl/ast_build.h"
|
|
|
|
#include "isl/set.h"
|
|
|
|
#include "isl/map.h"
|
|
|
|
#include "isl/aff.h"
|
|
|
|
|
2014-07-24 04:17:28 +08:00
|
|
|
#define DEBUG_TYPE "polly-ast"
|
|
|
|
|
2012-10-03 03:50:38 +08:00
|
|
|
using namespace llvm;
|
|
|
|
using namespace polly;
|
|
|
|
|
2014-07-24 04:17:28 +08:00
|
|
|
using IslAstUserPayload = IslAstInfo::IslAstUserPayload;
|
2014-04-22 11:30:19 +08:00
|
|
|
|
2013-05-07 15:30:56 +08:00
|
|
|
static cl::opt<bool> UseContext("polly-ast-use-context",
|
|
|
|
cl::desc("Use context"), cl::Hidden,
|
2013-05-07 15:31:10 +08:00
|
|
|
cl::init(false), cl::ZeroOrMore,
|
|
|
|
cl::cat(PollyCategory));
|
2012-10-03 03:50:38 +08:00
|
|
|
|
2013-05-07 15:30:56 +08:00
|
|
|
static cl::opt<bool> DetectParallel("polly-ast-detect-parallel",
|
|
|
|
cl::desc("Detect parallelism"), cl::Hidden,
|
2013-05-07 15:31:10 +08:00
|
|
|
cl::init(false), cl::ZeroOrMore,
|
|
|
|
cl::cat(PollyCategory));
|
2012-12-13 14:24:06 +08:00
|
|
|
|
2012-10-03 03:50:38 +08:00
|
|
|
namespace polly {
|
|
|
|
class IslAst {
|
|
|
|
public:
|
2012-12-13 14:24:06 +08:00
|
|
|
IslAst(Scop *Scop, Dependences &D);
|
2012-10-03 03:50:38 +08:00
|
|
|
|
|
|
|
~IslAst();
|
|
|
|
|
|
|
|
/// Print a source code representation of the program.
|
|
|
|
void pprint(llvm::raw_ostream &OS);
|
|
|
|
|
|
|
|
__isl_give isl_ast_node *getAst();
|
2013-10-31 19:50:52 +08:00
|
|
|
|
|
|
|
/// @brief Get the run-time conditions for the Scop.
|
2013-10-30 05:05:49 +08:00
|
|
|
__isl_give isl_ast_expr *getRunCondition();
|
2012-10-03 03:50:38 +08:00
|
|
|
|
|
|
|
private:
|
|
|
|
Scop *S;
|
|
|
|
isl_ast_node *Root;
|
2013-10-30 05:05:49 +08:00
|
|
|
isl_ast_expr *RunCondition;
|
2012-10-03 03:50:38 +08:00
|
|
|
|
2014-07-24 04:17:28 +08:00
|
|
|
void buildRunCondition(__isl_keep isl_ast_build *Build);
|
2012-10-03 03:50:38 +08:00
|
|
|
};
|
|
|
|
} // End namespace polly.
|
|
|
|
|
2014-07-24 04:17:28 +08:00
|
|
|
/// @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);
|
|
|
|
}
|
|
|
|
|
2014-07-24 23:59:06 +08:00
|
|
|
/// @brief Temporary information used when building the ast.
|
2012-12-13 14:24:06 +08:00
|
|
|
struct AstBuildUserInfo {
|
2014-07-24 23:59:06 +08:00
|
|
|
/// @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.
|
2012-12-13 14:24:06 +08:00
|
|
|
Dependences *Deps;
|
|
|
|
|
2014-07-24 23:59:06 +08:00
|
|
|
/// @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;
|
2012-12-13 14:24:06 +08:00
|
|
|
};
|
|
|
|
|
2014-08-01 05:33:49 +08:00
|
|
|
/// @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) {
|
|
|
|
Printer = isl_printer_start_line(Printer);
|
|
|
|
Printer = isl_printer_print_str(Printer, str.c_str());
|
|
|
|
return isl_printer_end_line(Printer);
|
2012-12-13 14:24:06 +08:00
|
|
|
}
|
|
|
|
|
2014-08-01 16:17:19 +08:00
|
|
|
/// @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;
|
|
|
|
|
|
|
|
BrokenReductions = IslAstInfo::getBrokenReductions(Node);
|
|
|
|
if (!BrokenReductions || BrokenReductions->empty())
|
|
|
|
return "";
|
|
|
|
|
|
|
|
// Map each type of reduction to a comma separated list of the base addresses.
|
|
|
|
std::map<MemoryAccess::ReductionType, std::string> Clauses;
|
|
|
|
for (MemoryAccess *MA : *BrokenReductions)
|
|
|
|
if (MA->isWrite())
|
|
|
|
Clauses[MA->getReductionType()] +=
|
|
|
|
", " + MA->getBaseAddr()->getName().str();
|
|
|
|
|
|
|
|
// Now print the reductions sorted by type. Each type will cause a clause
|
|
|
|
// like: reduction (+ : sum0, sum1, sum2)
|
|
|
|
for (const auto &ReductionClause : Clauses) {
|
|
|
|
str += " reduction (";
|
|
|
|
str += MemoryAccess::getReductionOperatorStr(ReductionClause.first);
|
|
|
|
// Remove the first two symbols (", ") to make the output look pretty.
|
|
|
|
str += " : " + ReductionClause.second.substr(2) + ")";
|
|
|
|
}
|
|
|
|
|
|
|
|
return str;
|
|
|
|
}
|
|
|
|
|
2014-08-01 05:33:49 +08:00
|
|
|
/// @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 *) {
|
2012-12-13 14:24:06 +08:00
|
|
|
|
2014-08-01 16:17:19 +08:00
|
|
|
const std::string BrokenReductionsStr = getBrokenReductionsStr(Node);
|
|
|
|
const std::string SimdPragmaStr = "#pragma simd";
|
|
|
|
const std::string OmpPragmaStr = "#pragma omp parallel for";
|
2014-08-01 05:33:49 +08:00
|
|
|
|
2014-08-01 16:17:19 +08:00
|
|
|
if (IslAstInfo::isInnermostParallel(Node))
|
|
|
|
Printer = printLine(Printer, SimdPragmaStr + BrokenReductionsStr);
|
2014-08-01 05:33:49 +08:00
|
|
|
|
2014-08-01 16:17:19 +08:00
|
|
|
if (IslAstInfo::isOutermostParallel(Node))
|
|
|
|
Printer = printLine(Printer, OmpPragmaStr + BrokenReductionsStr);
|
2014-08-01 05:33:49 +08:00
|
|
|
|
|
|
|
return isl_ast_node_for_print(Node, Printer, Options);
|
2012-12-13 14:24:06 +08:00
|
|
|
}
|
|
|
|
|
2014-07-15 08:00:35 +08:00
|
|
|
/// @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,
|
2014-08-01 16:17:19 +08:00
|
|
|
IslAstUserPayload *NodeInfo) {
|
2014-07-15 08:00:35 +08:00
|
|
|
if (!D->hasValidDependences())
|
|
|
|
return false;
|
|
|
|
|
2014-07-28 11:46:28 +08:00
|
|
|
isl_union_map *Schedule = isl_ast_build_get_schedule(Build);
|
2014-07-15 08:00:35 +08:00
|
|
|
isl_union_map *Deps = D->getDependences(
|
|
|
|
Dependences::TYPE_RAW | Dependences::TYPE_WAW | Dependences::TYPE_WAR);
|
2014-07-28 11:46:28 +08:00
|
|
|
if (!D->isParallel(Schedule, Deps) && !isl_union_map_free(Schedule))
|
2014-07-15 08:00:35 +08:00
|
|
|
return false;
|
|
|
|
|
2014-07-15 08:58:57 +08:00
|
|
|
isl_union_map *RedDeps = D->getDependences(Dependences::TYPE_TC_RED);
|
2014-07-28 11:46:28 +08:00
|
|
|
if (!D->isParallel(Schedule, RedDeps))
|
2014-08-01 16:17:19 +08:00
|
|
|
NodeInfo->IsReductionParallel = true;
|
|
|
|
|
|
|
|
if (!NodeInfo->IsReductionParallel && !isl_union_map_free(Schedule))
|
|
|
|
return true;
|
|
|
|
|
|
|
|
// Annotate reduction parallel nodes with the memory accesses which caused the
|
|
|
|
// reduction dependences parallel execution of the node conflicts with.
|
|
|
|
for (const auto &MaRedPair : D->getReductionDependences()) {
|
|
|
|
if (!MaRedPair.second)
|
|
|
|
continue;
|
|
|
|
RedDeps = isl_union_map_from_map(isl_map_copy(MaRedPair.second));
|
|
|
|
if (!D->isParallel(Schedule, RedDeps))
|
|
|
|
NodeInfo->BrokenReductions.insert(MaRedPair.first);
|
|
|
|
}
|
2014-07-15 08:00:35 +08:00
|
|
|
|
2014-07-28 11:46:28 +08:00
|
|
|
isl_union_map_free(Schedule);
|
2014-07-15 08:00:35 +08:00
|
|
|
return true;
|
|
|
|
}
|
|
|
|
|
2012-12-13 14:24:06 +08:00
|
|
|
// This method is executed before the construction of a for node. It creates
|
|
|
|
// an isl_id that is used to annotate the subsequently generated ast for nodes.
|
|
|
|
//
|
|
|
|
// In this function we also run the following analyses:
|
|
|
|
//
|
|
|
|
// - Detection of openmp parallel loops
|
|
|
|
//
|
2013-05-07 15:30:56 +08:00
|
|
|
static __isl_give isl_id *astBuildBeforeFor(__isl_keep isl_ast_build *Build,
|
|
|
|
void *User) {
|
2014-07-24 04:17:28 +08:00
|
|
|
AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
|
2014-08-01 05:34:32 +08:00
|
|
|
IslAstUserPayload *Payload = new IslAstUserPayload();
|
|
|
|
isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", Payload);
|
2014-07-24 04:17:28 +08:00
|
|
|
Id = isl_id_set_free_user(Id, freeIslAstUserPayload);
|
2014-07-24 23:59:06 +08:00
|
|
|
BuildInfo->LastForNodeId = Id;
|
2012-12-13 14:24:06 +08:00
|
|
|
|
2014-08-01 05:34:32 +08:00
|
|
|
// Test for parallelism only if we are not already inside a parallel loop
|
|
|
|
if (!BuildInfo->InParallelFor)
|
|
|
|
BuildInfo->InParallelFor = Payload->IsOutermostParallel =
|
2014-08-01 16:17:19 +08:00
|
|
|
astScheduleDimIsParallel(Build, BuildInfo->Deps, Payload);
|
2012-12-13 14:24:06 +08:00
|
|
|
|
|
|
|
return Id;
|
|
|
|
}
|
|
|
|
|
|
|
|
// This method is executed after the construction of a for node.
|
|
|
|
//
|
|
|
|
// It performs the following actions:
|
|
|
|
//
|
|
|
|
// - Reset the 'InParallelFor' flag, as soon as we leave a for node,
|
|
|
|
// that is marked as openmp parallel.
|
|
|
|
//
|
|
|
|
static __isl_give isl_ast_node *
|
2013-02-06 02:01:29 +08:00
|
|
|
astBuildAfterFor(__isl_take isl_ast_node *Node, __isl_keep isl_ast_build *Build,
|
|
|
|
void *User) {
|
2012-12-14 00:52:41 +08:00
|
|
|
isl_id *Id = isl_ast_node_get_annotation(Node);
|
2014-08-01 05:34:32 +08:00
|
|
|
assert(Id && "Post order visit assumes annotated for nodes");
|
|
|
|
IslAstUserPayload *Payload = (IslAstUserPayload *)isl_id_get_user(Id);
|
|
|
|
assert(Payload && "Post order visit assumes annotated for nodes");
|
2014-08-01 05:33:49 +08:00
|
|
|
|
2014-08-01 05:34:32 +08:00
|
|
|
AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
|
|
|
|
assert(!Payload->Build && "Build environment already set");
|
|
|
|
Payload->Build = isl_ast_build_copy(Build);
|
|
|
|
Payload->IsInnermost = (Id == BuildInfo->LastForNodeId);
|
|
|
|
|
|
|
|
// Innermost loops that are surrounded by parallel loops have not yet been
|
|
|
|
// tested for parallelism. Test them here to ensure we check all innermost
|
|
|
|
// loops for parallelism.
|
2014-08-01 16:17:19 +08:00
|
|
|
if (Payload->IsInnermost && BuildInfo->InParallelFor) {
|
2014-08-01 05:34:32 +08:00
|
|
|
if (Payload->IsOutermostParallel)
|
|
|
|
Payload->IsInnermostParallel = true;
|
|
|
|
else
|
2014-08-01 16:17:19 +08:00
|
|
|
Payload->IsInnermostParallel =
|
|
|
|
astScheduleDimIsParallel(Build, BuildInfo->Deps, Payload);
|
2014-09-10 01:03:54 +08:00
|
|
|
}
|
|
|
|
if (Payload->IsOutermostParallel)
|
2014-08-01 05:34:32 +08:00
|
|
|
BuildInfo->InParallelFor = false;
|
2012-12-13 14:24:06 +08:00
|
|
|
|
2012-12-18 15:46:13 +08:00
|
|
|
isl_id_free(Id);
|
2012-12-13 14:24:06 +08:00
|
|
|
return Node;
|
|
|
|
}
|
|
|
|
|
2013-05-07 15:30:56 +08:00
|
|
|
static __isl_give isl_ast_node *AtEachDomain(__isl_take isl_ast_node *Node,
|
2014-07-24 04:17:28 +08:00
|
|
|
__isl_keep isl_ast_build *Build,
|
2013-05-07 15:30:56 +08:00
|
|
|
void *User) {
|
2014-07-29 16:59:56 +08:00
|
|
|
assert(!isl_ast_node_get_annotation(Node) && "Node already annotated");
|
2014-08-01 05:34:32 +08:00
|
|
|
|
|
|
|
IslAstUserPayload *Payload = new IslAstUserPayload();
|
|
|
|
isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", Payload);
|
2014-07-29 16:59:56 +08:00
|
|
|
Id = isl_id_set_free_user(Id, freeIslAstUserPayload);
|
2012-10-03 03:50:38 +08:00
|
|
|
|
2014-08-01 05:34:32 +08:00
|
|
|
Payload->Build = isl_ast_build_copy(Build);
|
2012-10-03 03:50:38 +08:00
|
|
|
|
2012-12-18 15:46:13 +08:00
|
|
|
return isl_ast_node_set_annotation(Node, Id);
|
2012-10-03 03:50:38 +08:00
|
|
|
}
|
|
|
|
|
2014-07-24 04:17:28 +08:00
|
|
|
void IslAst::buildRunCondition(__isl_keep isl_ast_build *Build) {
|
2013-10-31 19:50:52 +08:00
|
|
|
// The conditions that need to be checked at run-time for this scop are
|
|
|
|
// available as an isl_set in the AssumedContext. We generate code for this
|
|
|
|
// check as follows. First, we generate an isl_pw_aff that is 1, if a certain
|
|
|
|
// combination of parameter values fulfills the conditions in the assumed
|
|
|
|
// context, and that is 0 otherwise. We then translate this isl_pw_aff into
|
|
|
|
// an isl_ast_expr. At run-time this expression can be evaluated and the
|
|
|
|
// optimized scop can be executed conditionally according to the result of the
|
|
|
|
// run-time check.
|
|
|
|
|
2013-10-30 05:05:49 +08:00
|
|
|
isl_aff *Zero =
|
|
|
|
isl_aff_zero_on_domain(isl_local_space_from_space(S->getParamSpace()));
|
|
|
|
isl_aff *One =
|
|
|
|
isl_aff_zero_on_domain(isl_local_space_from_space(S->getParamSpace()));
|
|
|
|
|
|
|
|
One = isl_aff_add_constant_si(One, 1);
|
|
|
|
|
|
|
|
isl_pw_aff *PwZero = isl_pw_aff_from_aff(Zero);
|
|
|
|
isl_pw_aff *PwOne = isl_pw_aff_from_aff(One);
|
|
|
|
|
|
|
|
PwOne = isl_pw_aff_intersect_domain(PwOne, S->getAssumedContext());
|
|
|
|
PwZero = isl_pw_aff_intersect_domain(
|
|
|
|
PwZero, isl_set_complement(S->getAssumedContext()));
|
|
|
|
|
2014-07-03 01:47:48 +08:00
|
|
|
isl_pw_aff *Cond = isl_pw_aff_union_max(PwOne, PwZero);
|
2013-10-30 05:05:49 +08:00
|
|
|
|
2014-07-24 04:17:28 +08:00
|
|
|
RunCondition = isl_ast_build_expr_from_pw_aff(Build, Cond);
|
2013-10-30 05:05:49 +08:00
|
|
|
}
|
|
|
|
|
2012-12-13 14:24:06 +08:00
|
|
|
IslAst::IslAst(Scop *Scop, Dependences &D) : S(Scop) {
|
2012-10-03 03:50:38 +08:00
|
|
|
isl_ctx *Ctx = S->getIslCtx();
|
|
|
|
isl_options_set_ast_build_atomic_upper_bound(Ctx, true);
|
2014-07-24 04:17:28 +08:00
|
|
|
isl_ast_build *Build;
|
|
|
|
AstBuildUserInfo BuildInfo;
|
2012-10-03 03:50:38 +08:00
|
|
|
|
|
|
|
if (UseContext)
|
2014-07-24 04:17:28 +08:00
|
|
|
Build = isl_ast_build_from_context(S->getContext());
|
2012-10-03 03:50:38 +08:00
|
|
|
else
|
2014-07-24 04:17:28 +08:00
|
|
|
Build = isl_ast_build_from_context(isl_set_universe(S->getParamSpace()));
|
2012-10-03 03:50:38 +08:00
|
|
|
|
2014-07-24 04:17:28 +08:00
|
|
|
Build = isl_ast_build_set_at_each_domain(Build, AtEachDomain, nullptr);
|
2012-10-03 03:50:38 +08:00
|
|
|
|
2014-06-28 16:59:38 +08:00
|
|
|
isl_union_map *Schedule =
|
|
|
|
isl_union_map_intersect_domain(S->getSchedule(), S->getDomains());
|
2012-10-03 03:50:38 +08:00
|
|
|
|
2012-12-18 15:46:13 +08:00
|
|
|
if (DetectParallel || PollyVectorizerChoice != VECTORIZER_NONE) {
|
2012-12-13 14:24:06 +08:00
|
|
|
BuildInfo.Deps = &D;
|
|
|
|
BuildInfo.InParallelFor = 0;
|
|
|
|
|
2014-07-24 04:17:28 +08:00
|
|
|
Build = isl_ast_build_set_before_each_for(Build, &astBuildBeforeFor,
|
|
|
|
&BuildInfo);
|
|
|
|
Build =
|
|
|
|
isl_ast_build_set_after_each_for(Build, &astBuildAfterFor, &BuildInfo);
|
2012-12-13 14:24:06 +08:00
|
|
|
}
|
|
|
|
|
2014-07-24 04:17:28 +08:00
|
|
|
buildRunCondition(Build);
|
2013-10-30 05:05:49 +08:00
|
|
|
|
2014-07-24 04:17:28 +08:00
|
|
|
Root = isl_ast_build_ast_from_schedule(Build, Schedule);
|
2012-10-03 03:50:38 +08:00
|
|
|
|
2014-07-24 04:17:28 +08:00
|
|
|
isl_ast_build_free(Build);
|
2012-10-03 03:50:38 +08:00
|
|
|
}
|
|
|
|
|
2013-10-30 05:05:49 +08:00
|
|
|
IslAst::~IslAst() {
|
|
|
|
isl_ast_node_free(Root);
|
|
|
|
isl_ast_expr_free(RunCondition);
|
|
|
|
}
|
2012-10-03 03:50:38 +08:00
|
|
|
|
2013-02-06 02:01:29 +08:00
|
|
|
__isl_give isl_ast_node *IslAst::getAst() { return isl_ast_node_copy(Root); }
|
2013-10-30 05:05:49 +08:00
|
|
|
__isl_give isl_ast_expr *IslAst::getRunCondition() {
|
|
|
|
return isl_ast_expr_copy(RunCondition);
|
|
|
|
}
|
2012-10-03 03:50:38 +08:00
|
|
|
|
|
|
|
void IslAstInfo::releaseMemory() {
|
|
|
|
if (Ast) {
|
|
|
|
delete Ast;
|
|
|
|
Ast = 0;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
bool IslAstInfo::runOnScop(Scop &Scop) {
|
|
|
|
if (Ast)
|
|
|
|
delete Ast;
|
|
|
|
|
|
|
|
S = &Scop;
|
|
|
|
|
2012-12-13 14:24:06 +08:00
|
|
|
Dependences &D = getAnalysis<Dependences>();
|
|
|
|
|
|
|
|
Ast = new IslAst(&Scop, D);
|
2012-10-03 03:50:38 +08:00
|
|
|
|
2014-07-24 02:14:43 +08:00
|
|
|
DEBUG(printScop(dbgs()));
|
2012-10-03 03:50:38 +08:00
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
2014-07-24 02:14:43 +08:00
|
|
|
__isl_give isl_ast_node *IslAstInfo::getAst() const { return Ast->getAst(); }
|
|
|
|
__isl_give isl_ast_expr *IslAstInfo::getRunCondition() const {
|
2013-11-17 11:18:25 +08:00
|
|
|
return Ast->getRunCondition();
|
|
|
|
}
|
2012-10-03 03:50:38 +08:00
|
|
|
|
2014-07-18 00:11:28 +08:00
|
|
|
IslAstUserPayload *IslAstInfo::getNodePayload(__isl_keep isl_ast_node *Node) {
|
|
|
|
isl_id *Id = isl_ast_node_get_annotation(Node);
|
|
|
|
if (!Id)
|
|
|
|
return nullptr;
|
|
|
|
IslAstUserPayload *Payload = (IslAstUserPayload *)isl_id_get_user(Id);
|
|
|
|
isl_id_free(Id);
|
|
|
|
return Payload;
|
|
|
|
}
|
|
|
|
|
2014-08-01 05:33:49 +08:00
|
|
|
bool IslAstInfo::isInnermost(__isl_keep isl_ast_node *Node) {
|
|
|
|
IslAstUserPayload *Payload = getNodePayload(Node);
|
|
|
|
return Payload && Payload->IsInnermost;
|
|
|
|
}
|
|
|
|
|
2014-07-18 00:11:28 +08:00
|
|
|
bool IslAstInfo::isParallel(__isl_keep isl_ast_node *Node) {
|
2014-08-01 16:14:28 +08:00
|
|
|
return IslAstInfo::isInnermostParallel(Node) ||
|
|
|
|
IslAstInfo::isOutermostParallel(Node);
|
2014-07-18 00:11:28 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
bool IslAstInfo::isInnermostParallel(__isl_keep isl_ast_node *Node) {
|
|
|
|
IslAstUserPayload *Payload = getNodePayload(Node);
|
2014-08-01 16:14:28 +08:00
|
|
|
return Payload && Payload->IsInnermostParallel;
|
2014-07-18 00:11:28 +08:00
|
|
|
}
|
|
|
|
|
2014-08-01 05:34:32 +08:00
|
|
|
bool IslAstInfo::isOutermostParallel(__isl_keep isl_ast_node *Node) {
|
2014-07-18 00:11:28 +08:00
|
|
|
IslAstUserPayload *Payload = getNodePayload(Node);
|
2014-08-01 16:14:28 +08:00
|
|
|
return Payload && Payload->IsOutermostParallel;
|
2014-07-18 00:11:28 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
bool IslAstInfo::isReductionParallel(__isl_keep isl_ast_node *Node) {
|
|
|
|
IslAstUserPayload *Payload = getNodePayload(Node);
|
|
|
|
return Payload && Payload->IsReductionParallel;
|
|
|
|
}
|
|
|
|
|
2014-07-24 04:17:28 +08:00
|
|
|
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;
|
|
|
|
}
|
|
|
|
|
2014-08-01 16:17:19 +08:00
|
|
|
IslAstInfo::MemoryAccessSet *
|
|
|
|
IslAstInfo::getBrokenReductions(__isl_keep isl_ast_node *Node) {
|
|
|
|
IslAstUserPayload *Payload = getNodePayload(Node);
|
|
|
|
return Payload ? &Payload->BrokenReductions : nullptr;
|
|
|
|
}
|
|
|
|
|
2014-08-03 09:51:59 +08:00
|
|
|
isl_ast_build *IslAstInfo::getBuild(__isl_keep isl_ast_node *Node) {
|
|
|
|
IslAstUserPayload *Payload = getNodePayload(Node);
|
|
|
|
return Payload ? Payload->Build : nullptr;
|
|
|
|
}
|
|
|
|
|
2014-07-24 02:14:43 +08:00
|
|
|
void IslAstInfo::printScop(raw_ostream &OS) const {
|
|
|
|
isl_ast_print_options *Options;
|
|
|
|
isl_ast_node *RootNode = getAst();
|
|
|
|
isl_ast_expr *RunCondition = getRunCondition();
|
|
|
|
char *RtCStr, *AstStr;
|
|
|
|
|
|
|
|
Scop &S = getCurScop();
|
|
|
|
Options = isl_ast_print_options_alloc(S.getIslCtx());
|
2014-08-01 05:33:49 +08:00
|
|
|
Options = isl_ast_print_options_set_print_for(Options, cbPrintFor, nullptr);
|
2014-07-24 02:14:43 +08:00
|
|
|
|
|
|
|
isl_printer *P = isl_printer_to_str(S.getIslCtx());
|
|
|
|
P = isl_printer_print_ast_expr(P, RunCondition);
|
|
|
|
RtCStr = isl_printer_get_str(P);
|
|
|
|
P = isl_printer_flush(P);
|
|
|
|
P = isl_printer_indent(P, 4);
|
|
|
|
P = isl_printer_set_output_format(P, ISL_FORMAT_C);
|
|
|
|
P = isl_ast_node_print(RootNode, P, Options);
|
|
|
|
AstStr = isl_printer_get_str(P);
|
|
|
|
|
|
|
|
Function *F = S.getRegion().getEntry()->getParent();
|
|
|
|
isl_union_map *Schedule =
|
|
|
|
isl_union_map_intersect_domain(S.getSchedule(), S.getDomains());
|
|
|
|
|
|
|
|
OS << ":: isl ast :: " << F->getName() << " :: " << S.getRegion().getNameStr()
|
|
|
|
<< "\n";
|
|
|
|
DEBUG(dbgs() << S.getContextStr() << "\n"; isl_union_map_dump(Schedule));
|
|
|
|
OS << "\nif (" << RtCStr << ")\n\n";
|
|
|
|
OS << AstStr << "\n";
|
|
|
|
OS << "else\n";
|
|
|
|
OS << " { /* original code */ }\n\n";
|
|
|
|
|
|
|
|
isl_ast_expr_free(RunCondition);
|
|
|
|
isl_union_map_free(Schedule);
|
|
|
|
isl_ast_node_free(RootNode);
|
|
|
|
isl_printer_free(P);
|
|
|
|
}
|
|
|
|
|
2012-10-03 03:50:38 +08:00
|
|
|
void IslAstInfo::getAnalysisUsage(AnalysisUsage &AU) const {
|
|
|
|
// Get the Common analysis usage of ScopPasses.
|
|
|
|
ScopPass::getAnalysisUsage(AU);
|
|
|
|
AU.addRequired<ScopInfo>();
|
2012-12-13 14:24:06 +08:00
|
|
|
AU.addRequired<Dependences>();
|
2012-10-03 03:50:38 +08:00
|
|
|
}
|
2013-02-22 16:07:06 +08:00
|
|
|
|
2012-10-03 03:50:38 +08:00
|
|
|
char IslAstInfo::ID = 0;
|
|
|
|
|
2013-02-22 16:07:06 +08:00
|
|
|
Pass *polly::createIslAstInfoPass() { return new IslAstInfo(); }
|
|
|
|
|
2012-10-03 03:50:38 +08:00
|
|
|
INITIALIZE_PASS_BEGIN(IslAstInfo, "polly-ast",
|
2013-03-23 09:05:07 +08:00
|
|
|
"Polly - Generate an AST of the SCoP (isl)", false,
|
|
|
|
false);
|
2013-02-22 16:07:06 +08:00
|
|
|
INITIALIZE_PASS_DEPENDENCY(ScopInfo);
|
|
|
|
INITIALIZE_PASS_DEPENDENCY(Dependences);
|
2012-10-03 03:50:38 +08:00
|
|
|
INITIALIZE_PASS_END(IslAstInfo, "polly-ast",
|
2013-03-23 09:05:07 +08:00
|
|
|
"Polly - Generate an AST from the SCoP (isl)", false, false)
|