maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/polly/lib/CodeGen/IslAst.cpp

851 lines
30 KiB
C++
Raw Blame History

//===- 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 an 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);
// }
//
// An in-depth discussion of our AST generation approach can be found in:
//
// Polyhedral AST generation is more than scanning polyhedra
// Tobias Grosser, Sven Verdoolaege, Albert Cohen
// ACM Transactions on Programming Languages and Systems (TOPLAS),
// 37(4), July 2015
// http://www.grosser.es/#pub-polyhedral-AST-generation
//
//===----------------------------------------------------------------------===//
#include "polly/CodeGen/IslAst.h"
#include "polly/CodeGen/CodeGeneration.h"
#include "polly/DependenceInfo.h"
#include "polly/LinkAllPasses.h"
#include "polly/Options.h"
#include "polly/ScopDetection.h"
#include "polly/ScopInfo.h"
#include "polly/ScopPass.h"
#include "polly/Support/GICHelper.h"
#include "llvm/ADT/Statistic.h"
#include "llvm/IR/Function.h"
#include "llvm/Pass.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/Debug.h"
#include "llvm/Support/raw_ostream.h"
#include "isl/aff.h"
#include "isl/ast.h"
#include "isl/ast_build.h"
#include "isl/id.h"
#include "isl/isl-noexceptions.h"
#include "isl/map.h"
#include "isl/printer.h"
#include "isl/schedule.h"
#include "isl/set.h"
#include "isl/union_map.h"
#include "isl/val.h"
#include <cassert>
#include <cstdlib>
#include <cstring>
#include <map>
#include <string>
#include <utility>
#define DEBUG_TYPE "polly-ast"
using namespace llvm;
using namespace polly;
using IslAstUserPayload = IslAstInfo::IslAstUserPayload;
static cl::opt<bool>
PollyParallel("polly-parallel",
cl::desc("Generate thread parallel code (isl codegen only)"),
cl::init(false), cl::ZeroOrMore, cl::cat(PollyCategory));
static cl::opt<bool> PrintAccesses("polly-ast-print-accesses",
cl::desc("Print memory access functions"),
cl::init(false), cl::ZeroOrMore,
cl::cat(PollyCategory));
static cl::opt<bool> PollyParallelForce(
"polly-parallel-force",
cl::desc(
"Force generation of thread parallel code ignoring any cost model"),
cl::init(false), cl::ZeroOrMore, cl::cat(PollyCategory));
static cl::opt<bool> UseContext("polly-ast-use-context",
cl::desc("Use context"), cl::Hidden,
cl::init(true), cl::ZeroOrMore,
cl::cat(PollyCategory));
static cl::opt<bool> DetectParallel("polly-ast-detect-parallel",
cl::desc("Detect parallelism"), cl::Hidden,
cl::init(false), cl::ZeroOrMore,
cl::cat(PollyCategory));
STATISTIC(ScopsProcessed, "Number of SCoPs processed");
STATISTIC(ScopsBeneficial, "Number of beneficial SCoPs");
STATISTIC(BeneficialAffineLoops, "Number of beneficial affine loops");
STATISTIC(BeneficialBoxedLoops, "Number of beneficial boxed loops");
STATISTIC(NumForLoops, "Number of for-loops");
STATISTIC(NumParallel, "Number of parallel for-loops");
STATISTIC(NumInnermostParallel, "Number of innermost parallel for-loops");
STATISTIC(NumOutermostParallel, "Number of outermost parallel for-loops");
STATISTIC(NumReductionParallel, "Number of reduction-parallel for-loops");
STATISTIC(NumExecutedInParallel, "Number of for-loops executed in parallel");
STATISTIC(NumIfConditions, "Number of if-conditions");
namespace polly {
/// Temporary information used when building the ast.
struct AstBuildUserInfo {
/// Construct and initialize the helper struct for AST creation.
AstBuildUserInfo() = default;
/// The dependence information used for the parallelism check.
const Dependences *Deps = nullptr;
/// Flag to indicate that we are inside a parallel for node.
bool InParallelFor = false;
/// Flag to indicate that we are inside an SIMD node.
bool InSIMD = false;
/// The last iterator id created for the current SCoP.
isl_id *LastForNodeId = nullptr;
};
} // namespace polly
/// 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);
}
/// 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,
__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);
}
/// 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->getScopArrayInfo()->getName();
// 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;
}
/// 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::getMinimalDependenceDistance(Node);
const std::string BrokenReductionsStr = getBrokenReductionsStr(Node);
const std::string KnownParallelStr = "#pragma known-parallel";
const std::string DepDisPragmaStr = "#pragma minimal 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::isExecutedInParallel(Node))
Printer = printLine(Printer, OmpPragmaStr);
else if (IslAstInfo::isOutermostParallel(Node))
Printer = printLine(Printer, KnownParallelStr + BrokenReductionsStr);
isl_pw_aff_free(DD);
return isl_ast_node_for_print(Node, Printer, Options);
}
/// 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 reversal. 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,
const 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)
.release();
if (!D->isParallel(Schedule, Deps)) {
isl_union_map *DepsAll =
D->getDependences(Dependences::TYPE_RAW | Dependences::TYPE_WAW |
Dependences::TYPE_WAR | Dependences::TYPE_TC_RED)
.release();
isl_pw_aff *MinimalDependenceDistance = nullptr;
D->isParallel(Schedule, DepsAll, &MinimalDependenceDistance);
NodeInfo->MinimalDependenceDistance =
isl::manage(MinimalDependenceDistance);
isl_union_map_free(Schedule);
return false;
}
isl_union_map *RedDeps =
D->getDependences(Dependences::TYPE_TC_RED).release();
if (!D->isParallel(Schedule, RedDeps))
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);
}
isl_union_map_free(Schedule);
return true;
}
// 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
//
static __isl_give isl_id *astBuildBeforeFor(__isl_keep isl_ast_build *Build,
void *User) {
AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
IslAstUserPayload *Payload = new IslAstUserPayload();
isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", Payload);
Id = isl_id_set_free_user(Id, freeIslAstUserPayload);
BuildInfo->LastForNodeId = Id;
Payload->IsParallel =
astScheduleDimIsParallel(Build, BuildInfo->Deps, Payload);
// Test for parallelism only if we are not already inside a parallel loop
if (!BuildInfo->InParallelFor && !BuildInfo->InSIMD)
BuildInfo->InParallelFor = Payload->IsOutermostParallel =
Payload->IsParallel;
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 *
astBuildAfterFor(__isl_take isl_ast_node *Node, __isl_keep isl_ast_build *Build,
void *User) {
isl_id *Id = isl_ast_node_get_annotation(Node);
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");
AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
assert(!Payload->Build && "Build environment already set");
Payload->Build = isl_ast_build_copy(Build);
Payload->IsInnermost = (Id == BuildInfo->LastForNodeId);
Payload->IsInnermostParallel =
Payload->IsInnermost && (BuildInfo->InSIMD || Payload->IsParallel);
if (Payload->IsOutermostParallel)
BuildInfo->InParallelFor = false;
isl_id_free(Id);
return Node;
}
static isl_stat astBuildBeforeMark(__isl_keep isl_id *MarkId,
__isl_keep isl_ast_build *Build,
void *User) {
if (!MarkId)
return isl_stat_error;
AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
if (strcmp(isl_id_get_name(MarkId), "SIMD") == 0)
BuildInfo->InSIMD = true;
return isl_stat_ok;
}
static __isl_give isl_ast_node *
astBuildAfterMark(__isl_take isl_ast_node *Node,
__isl_keep isl_ast_build *Build, void *User) {
assert(isl_ast_node_get_type(Node) == isl_ast_node_mark);
AstBuildUserInfo *BuildInfo = (AstBuildUserInfo *)User;
auto *Id = isl_ast_node_mark_get_id(Node);
if (strcmp(isl_id_get_name(Id), "SIMD") == 0)
BuildInfo->InSIMD = false;
isl_id_free(Id);
return Node;
}
static __isl_give isl_ast_node *AtEachDomain(__isl_take isl_ast_node *Node,
__isl_keep isl_ast_build *Build,
void *User) {
assert(!isl_ast_node_get_annotation(Node) && "Node already annotated");
IslAstUserPayload *Payload = new IslAstUserPayload();
isl_id *Id = isl_id_alloc(isl_ast_build_get_ctx(Build), "", Payload);
Id = isl_id_set_free_user(Id, freeIslAstUserPayload);
Payload->Build = isl_ast_build_copy(Build);
return isl_ast_node_set_annotation(Node, Id);
}
// Build alias check condition given a pair of minimal/maximal access.
static isl::ast_expr buildCondition(Scop &S, isl::ast_build Build,
const Scop::MinMaxAccessTy *It0,
const Scop::MinMaxAccessTy *It1) {
isl::pw_multi_aff AFirst = It0->first;
isl::pw_multi_aff ASecond = It0->second;
isl::pw_multi_aff BFirst = It1->first;
isl::pw_multi_aff BSecond = It1->second;
isl::id Left = AFirst.get_tuple_id(isl::dim::set);
isl::id Right = BFirst.get_tuple_id(isl::dim::set);
isl::ast_expr True =
isl::ast_expr::from_val(isl::val::int_from_ui(Build.get_ctx(), 1));
isl::ast_expr False =
isl::ast_expr::from_val(isl::val::int_from_ui(Build.get_ctx(), 0));
const ScopArrayInfo *BaseLeft =
ScopArrayInfo::getFromId(Left)->getBasePtrOriginSAI();
const ScopArrayInfo *BaseRight =
ScopArrayInfo::getFromId(Right)->getBasePtrOriginSAI();
if (BaseLeft && BaseLeft == BaseRight)
return True;
isl::set Params = S.getContext();
isl::ast_expr NonAliasGroup, MinExpr, MaxExpr;
// In the following, we first check if any accesses will be empty under
// the execution context of the scop and do not code generate them if this
// is the case as isl will fail to derive valid AST expressions for such
// accesses.
if (!AFirst.intersect_params(Params).domain().is_empty() &&
!BSecond.intersect_params(Params).domain().is_empty()) {
MinExpr = Build.access_from(AFirst).address_of();
MaxExpr = Build.access_from(BSecond).address_of();
NonAliasGroup = MaxExpr.le(MinExpr);
}
if (!BFirst.intersect_params(Params).domain().is_empty() &&
!ASecond.intersect_params(Params).domain().is_empty()) {
MinExpr = Build.access_from(BFirst).address_of();
MaxExpr = Build.access_from(ASecond).address_of();
isl::ast_expr Result = MaxExpr.le(MinExpr);
if (!NonAliasGroup.is_null())
NonAliasGroup = isl::manage(
isl_ast_expr_or(NonAliasGroup.release(), Result.release()));
else
NonAliasGroup = Result;
}
if (NonAliasGroup.is_null())
NonAliasGroup = True;
return NonAliasGroup;
}
__isl_give isl_ast_expr *
IslAst::buildRunCondition(Scop &S, __isl_keep isl_ast_build *Build) {
isl_ast_expr *RunCondition;
// The conditions that need to be checked at run-time for this scop are
// available as an isl_set in the runtime check context from which we can
// directly derive a run-time condition.
auto *PosCond =
isl_ast_build_expr_from_set(Build, S.getAssumedContext().release());
if (S.hasTrivialInvalidContext()) {
RunCondition = PosCond;
} else {
auto *ZeroV = isl_val_zero(isl_ast_build_get_ctx(Build));
auto *NegCond =
isl_ast_build_expr_from_set(Build, S.getInvalidContext().release());
auto *NotNegCond = isl_ast_expr_eq(isl_ast_expr_from_val(ZeroV), NegCond);
RunCondition = isl_ast_expr_and(PosCond, NotNegCond);
}
// Create the alias checks from the minimal/maximal accesses in each alias
// group which consists of read only and non read only (read write) accesses.
// This operation is by construction quadratic in the read-write pointers and
// linear in the read only pointers in each alias group.
for (const Scop::MinMaxVectorPairTy &MinMaxAccessPair : S.getAliasGroups()) {
auto &MinMaxReadWrite = MinMaxAccessPair.first;
auto &MinMaxReadOnly = MinMaxAccessPair.second;
auto RWAccEnd = MinMaxReadWrite.end();
for (auto RWAccIt0 = MinMaxReadWrite.begin(); RWAccIt0 != RWAccEnd;
++RWAccIt0) {
for (auto RWAccIt1 = RWAccIt0 + 1; RWAccIt1 != RWAccEnd; ++RWAccIt1)
RunCondition = isl_ast_expr_and(
RunCondition,
buildCondition(S, isl::manage_copy(Build), RWAccIt0, RWAccIt1)
.release());
for (const Scop::MinMaxAccessTy &ROAccIt : MinMaxReadOnly)
RunCondition = isl_ast_expr_and(
RunCondition,
buildCondition(S, isl::manage_copy(Build), RWAccIt0, &ROAccIt)
.release());
}
}
return RunCondition;
}
/// Simple cost analysis for a given SCoP.
///
/// TODO: Improve this analysis and extract it to make it usable in other
/// places too.
/// In order to improve the cost model we could either keep track of
/// performed optimizations (e.g., tiling) or compute properties on the
/// original as well as optimized SCoP (e.g., #stride-one-accesses).
static bool benefitsFromPolly(Scop &Scop, bool PerformParallelTest) {
if (PollyProcessUnprofitable)
return true;
// Check if nothing interesting happened.
if (!PerformParallelTest && !Scop.isOptimized() &&
Scop.getAliasGroups().empty())
return false;
// The default assumption is that Polly improves the code.
return true;
}
/// Collect statistics for the syntax tree rooted at @p Ast.
static void walkAstForStatistics(__isl_keep isl_ast_node *Ast) {
assert(Ast);
isl_ast_node_foreach_descendant_top_down(
Ast,
[](__isl_keep isl_ast_node *Node, void *User) -> isl_bool {
switch (isl_ast_node_get_type(Node)) {
case isl_ast_node_for:
NumForLoops++;
if (IslAstInfo::isParallel(Node))
NumParallel++;
if (IslAstInfo::isInnermostParallel(Node))
NumInnermostParallel++;
if (IslAstInfo::isOutermostParallel(Node))
NumOutermostParallel++;
if (IslAstInfo::isReductionParallel(Node))
NumReductionParallel++;
if (IslAstInfo::isExecutedInParallel(Node))
NumExecutedInParallel++;
break;
case isl_ast_node_if:
NumIfConditions++;
break;
default:
break;
}
// Continue traversing subtrees.
return isl_bool_true;
},
nullptr);
}
IslAst::IslAst(Scop &Scop) : S(Scop), Ctx(Scop.getSharedIslCtx()) {}
IslAst::IslAst(IslAst &&O)
: S(O.S), Root(O.Root), RunCondition(O.RunCondition), Ctx(O.Ctx) {
O.Root = nullptr;
O.RunCondition = nullptr;
}
IslAst::~IslAst() {
isl_ast_node_free(Root);
isl_ast_expr_free(RunCondition);
}
void IslAst::init(const Dependences &D) {
bool PerformParallelTest = PollyParallel || DetectParallel ||
PollyVectorizerChoice != VECTORIZER_NONE;
// We can not perform the dependence analysis and, consequently,
// the parallel code generation in case the schedule tree contains
// extension nodes.
auto ScheduleTree = S.getScheduleTree();
PerformParallelTest =
PerformParallelTest && !S.containsExtensionNode(ScheduleTree);
// Skip AST and code generation if there was no benefit achieved.
if (!benefitsFromPolly(S, PerformParallelTest))
return;
auto ScopStats = S.getStatistics();
ScopsBeneficial++;
BeneficialAffineLoops += ScopStats.NumAffineLoops;
BeneficialBoxedLoops += ScopStats.NumBoxedLoops;
auto Ctx = S.getIslCtx();
isl_options_set_ast_build_atomic_upper_bound(Ctx.get(), true);
isl_options_set_ast_build_detect_min_max(Ctx.get(), true);
isl_ast_build *Build;
AstBuildUserInfo BuildInfo;
if (UseContext)
Build = isl_ast_build_from_context(S.getContext().release());
else
Build = isl_ast_build_from_context(
isl_set_universe(S.getParamSpace().release()));
Build = isl_ast_build_set_at_each_domain(Build, AtEachDomain, nullptr);
if (PerformParallelTest) {
BuildInfo.Deps = &D;
BuildInfo.InParallelFor = false;
BuildInfo.InSIMD = false;
Build = isl_ast_build_set_before_each_for(Build, &astBuildBeforeFor,
&BuildInfo);
Build =
isl_ast_build_set_after_each_for(Build, &astBuildAfterFor, &BuildInfo);
Build = isl_ast_build_set_before_each_mark(Build, &astBuildBeforeMark,
&BuildInfo);
Build = isl_ast_build_set_after_each_mark(Build, &astBuildAfterMark,
&BuildInfo);
}
RunCondition = buildRunCondition(S, Build);
Root = isl_ast_build_node_from_schedule(Build, S.getScheduleTree().release());
walkAstForStatistics(Root);
isl_ast_build_free(Build);
}
IslAst IslAst::create(Scop &Scop, const Dependences &D) {
IslAst Ast{Scop};
Ast.init(D);
return Ast;
}
__isl_give isl_ast_node *IslAst::getAst() { return isl_ast_node_copy(Root); }
__isl_give isl_ast_expr *IslAst::getRunCondition() {
return isl_ast_expr_copy(RunCondition);
}
__isl_give isl_ast_node *IslAstInfo::getAst() { return Ast.getAst(); }
__isl_give isl_ast_expr *IslAstInfo::getRunCondition() {
return Ast.getRunCondition();
}
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;
}
bool IslAstInfo::isInnermost(__isl_keep isl_ast_node *Node) {
IslAstUserPayload *Payload = getNodePayload(Node);
return Payload && Payload->IsInnermost;
}
bool IslAstInfo::isParallel(__isl_keep isl_ast_node *Node) {
return IslAstInfo::isInnermostParallel(Node) ||
IslAstInfo::isOutermostParallel(Node);
}
bool IslAstInfo::isInnermostParallel(__isl_keep isl_ast_node *Node) {
IslAstUserPayload *Payload = getNodePayload(Node);
return Payload && Payload->IsInnermostParallel;
}
bool IslAstInfo::isOutermostParallel(__isl_keep isl_ast_node *Node) {
IslAstUserPayload *Payload = getNodePayload(Node);
return Payload && Payload->IsOutermostParallel;
}
bool IslAstInfo::isReductionParallel(__isl_keep isl_ast_node *Node) {
IslAstUserPayload *Payload = getNodePayload(Node);
return Payload && Payload->IsReductionParallel;
}
bool IslAstInfo::isExecutedInParallel(__isl_keep isl_ast_node *Node) {
if (!PollyParallel)
return false;
// Do not parallelize innermost loops.
//
// Parallelizing innermost loops is often not profitable, especially if
// they have a low number of iterations.
//
// TODO: Decide this based on the number of loop iterations that will be
// executed. This can possibly require run-time checks, which again
// raises the question of both run-time check overhead and code size
// costs.
if (!PollyParallelForce && isInnermost(Node))
return false;
return isOutermostParallel(Node) && !isReductionParallel(Node);
}
__isl_give 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_give isl_pw_aff *
IslAstInfo::getMinimalDependenceDistance(__isl_keep isl_ast_node *Node) {
IslAstUserPayload *Payload = getNodePayload(Node);
return Payload ? Payload->MinimalDependenceDistance.copy() : 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);
return Payload ? Payload->Build : nullptr;
}
IslAstInfo IslAstAnalysis::run(Scop &S, ScopAnalysisManager &SAM,
ScopStandardAnalysisResults &SAR) {
return {S, SAM.getResult<DependenceAnalysis>(S, SAR).getDependences(
Dependences::AL_Statement)};
}
static __isl_give isl_printer *cbPrintUser(__isl_take isl_printer *P,
__isl_take isl_ast_print_options *O,
__isl_keep isl_ast_node *Node,
void *User) {
isl::ast_node AstNode = isl::manage_copy(Node);
isl::ast_expr NodeExpr = AstNode.user_get_expr();
isl::ast_expr CallExpr = NodeExpr.get_op_arg(0);
isl::id CallExprId = CallExpr.get_id();
ScopStmt *AccessStmt = (ScopStmt *)CallExprId.get_user();
P = isl_printer_start_line(P);
P = isl_printer_print_str(P, AccessStmt->getBaseName());
P = isl_printer_print_str(P, "(");
P = isl_printer_end_line(P);
P = isl_printer_indent(P, 2);
for (MemoryAccess *MemAcc : *AccessStmt) {
P = isl_printer_start_line(P);
if (MemAcc->isRead())
P = isl_printer_print_str(P, "/* read */ &");
else
P = isl_printer_print_str(P, "/* write */ ");
isl::ast_build Build = isl::manage_copy(IslAstInfo::getBuild(Node));
if (MemAcc->isAffine()) {
isl_pw_multi_aff *PwmaPtr =
MemAcc->applyScheduleToAccessRelation(Build.get_schedule()).release();
isl::pw_multi_aff Pwma = isl::manage(PwmaPtr);
isl::ast_expr AccessExpr = Build.access_from(Pwma);
P = isl_printer_print_ast_expr(P, AccessExpr.get());
} else {
P = isl_printer_print_str(
P, MemAcc->getLatestScopArrayInfo()->getName().c_str());
P = isl_printer_print_str(P, "[*]");
}
P = isl_printer_end_line(P);
}
P = isl_printer_indent(P, -2);
P = isl_printer_start_line(P);
P = isl_printer_print_str(P, ");");
P = isl_printer_end_line(P);
isl_ast_print_options_free(O);
return P;
}
void IslAstInfo::print(raw_ostream &OS) {
isl_ast_print_options *Options;
isl_ast_node *RootNode = Ast.getAst();
Function &F = S.getFunction();
OS << ":: isl ast :: " << F.getName() << " :: " << S.getNameStr() << "\n";
if (!RootNode) {
OS << ":: isl ast generation and code generation was skipped!\n\n";
OS << ":: This is either because no useful optimizations could be applied "
"(use -polly-process-unprofitable to enforce code generation) or "
"because earlier passes such as dependence analysis timed out (use "
"-polly-dependences-computeout=0 to set dependence analysis timeout "
"to infinity)\n\n";
return;
}
isl_ast_expr *RunCondition = Ast.getRunCondition();
char *RtCStr, *AstStr;
Options = isl_ast_print_options_alloc(S.getIslCtx().get());
if (PrintAccesses)
Options =
isl_ast_print_options_set_print_user(Options, cbPrintUser, nullptr);
Options = isl_ast_print_options_set_print_for(Options, cbPrintFor, nullptr);
isl_printer *P = isl_printer_to_str(S.getIslCtx().get());
P = isl_printer_set_output_format(P, ISL_FORMAT_C);
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_ast_node_print(RootNode, P, Options);
AstStr = isl_printer_get_str(P);
auto *Schedule = S.getScheduleTree().release();
LLVM_DEBUG({
dbgs() << S.getContextStr() << "\n";
dbgs() << stringFromIslObj(Schedule);
});
OS << "\nif (" << RtCStr << ")\n\n";
OS << AstStr << "\n";
OS << "else\n";
OS << " { /* original code */ }\n\n";
free(RtCStr);
free(AstStr);
isl_ast_expr_free(RunCondition);
isl_schedule_free(Schedule);
isl_ast_node_free(RootNode);
isl_printer_free(P);
}
AnalysisKey IslAstAnalysis::Key;
PreservedAnalyses IslAstPrinterPass::run(Scop &S, ScopAnalysisManager &SAM,
ScopStandardAnalysisResults &SAR,
SPMUpdater &U) {
auto &Ast = SAM.getResult<IslAstAnalysis>(S, SAR);
Ast.print(OS);
return PreservedAnalyses::all();
}
void IslAstInfoWrapperPass::releaseMemory() { Ast.reset(); }
bool IslAstInfoWrapperPass::runOnScop(Scop &Scop) {
// Skip SCoPs in case they're already handled by PPCGCodeGeneration.
if (Scop.isToBeSkipped())
return false;
ScopsProcessed++;
const Dependences &D =
getAnalysis<DependenceInfo>().getDependences(Dependences::AL_Statement);
if (D.getSharedIslCtx() != Scop.getSharedIslCtx()) {
LLVM_DEBUG(
dbgs() << "Got dependence analysis for different SCoP/isl_ctx\n");
Ast.reset();
return false;
}
Ast.reset(new IslAstInfo(Scop, D));
LLVM_DEBUG(printScop(dbgs(), Scop));
return false;
}
void IslAstInfoWrapperPass::getAnalysisUsage(AnalysisUsage &AU) const {
// Get the Common analysis usage of ScopPasses.
ScopPass::getAnalysisUsage(AU);
AU.addRequiredTransitive<ScopInfoRegionPass>();
AU.addRequired<DependenceInfo>();
AU.addPreserved<DependenceInfo>();
}
void IslAstInfoWrapperPass::printScop(raw_ostream &OS, Scop &S) const {
if (Ast)
Ast->print(OS);
}
char IslAstInfoWrapperPass::ID = 0;
Pass *polly::createIslAstInfoWrapperPassPass() {
return new IslAstInfoWrapperPass();
}
INITIALIZE_PASS_BEGIN(IslAstInfoWrapperPass, "polly-ast",
"Polly - Generate an AST of the SCoP (isl)", false,
false);
INITIALIZE_PASS_DEPENDENCY(ScopInfoRegionPass);
INITIALIZE_PASS_DEPENDENCY(DependenceInfo);
INITIALIZE_PASS_END(IslAstInfoWrapperPass, "polly-ast",
"Polly - Generate an AST from the SCoP (isl)", false, false)
Reference in New Issue View Git Blame Copy Permalink