[Polly] Annotate reduction dependences
Tobias Grosser
tobias at grosser.es
Wed Jul 30 12:47:07 PDT 2014
On 30/07/2014 19:05, Johannes Doerfert wrote:
> Here a second version.
>
> The reduction dependences stay in the Dependences class now and I added
> a print function in OpenMP style.
>
> --
>
> Johannes Doerfert
>
> Employee of Qualcomm Innovation Center, Inc.
>
> Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum,
> hosted by The Linux Foundation
>
> *From:*llvm-commits-bounces at cs.uiuc.edu
> [mailto:llvm-commits-bounces at cs.uiuc.edu] *On Behalf Of *Johannes Doerfert
> *Sent:* Thursday, July 17, 2014 5:22 PM
> *To:* llvm-commits at cs.uiuc.edu
> *Cc:* 'Sebastian Pop'; 'Tobias Grosser'; 'zino'
> *Subject:* [Polly] Annotate reduction dependences
>
> Hey,
>
> Attached is a patch to annotate reduction dependences caused by each
> memory access to this memory access.
>
> The reason we want to do that is code generation. At some point we need
> to determine which accesses need to be privatized in order to generate
> valid parallel/vectorized code.
>
> To find these accesses later we would need to iterate over all candidate
> memory accesses and perform the parallelism check again on the reduction
> dependences of the memory access.
>
> I’m open to different solutions but I think this is way is pretty save.
>
> Best regards,
>
> Johannes
>
> --
>
> Johannes Doerfert
>
> Employee of Qualcomm Innovation Center, Inc.
>
> Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum,
> hosted by The Linux Foundation
>
>
> 0001-Annotate-the-IslAst-with-broken-reductions.patch
>
>
> From 7b587583ffd8f38bee3021c73fcf3c6b5409e701 Mon Sep 17 00:00:00 2001
> From: Johannes Doerfert<jdoerfert at codeaurora.org>
> Date: Tue, 15 Jul 2014 16:13:27 -0700
> Subject: [PATCH] Annotate the IslAst with broken reductions
>
> + Print the types & base addresses of broken reductions for each "reduction
> parallel" marked loop.
> + 3 test cases
Can you add a little information into the commit message regarding why
this is useful and what it is doing. With this commit message, nobody
who does not know what you are doing gets this from a quick look.
> ---
> include/polly/CodeGen/IslAst.h | 14 +++-
> include/polly/Dependences.h | 18 ++++++
> lib/Analysis/Dependences.cpp | 38 +++++++++++
> lib/CodeGen/IslAst.cpp | 74 +++++++++++++++++++---
> .../reduction_clauses_multidimensional_access.ll | 68 ++++++++++++++++++++
> .../Ast/reduction_clauses_onedimensional_access.ll | 50 +++++++++++++++
> .../Ast/reduction_different_reduction_clauses.ll | 51 +++++++++++++++
> 7 files changed, 303 insertions(+), 10 deletions(-)
> create mode 100644 test/Isl/Ast/reduction_clauses_multidimensional_access.ll
> create mode 100644 test/Isl/Ast/reduction_clauses_onedimensional_access.ll
> create mode 100644 test/Isl/Ast/reduction_different_reduction_clauses.ll
>
> diff --git a/include/polly/CodeGen/IslAst.h b/include/polly/CodeGen/IslAst.h
> index 341d56c..1406fe4 100644
> --- a/include/polly/CodeGen/IslAst.h
> +++ b/include/polly/CodeGen/IslAst.h
> @@ -34,15 +34,19 @@ class raw_ostream;
> struct isl_ast_node;
> struct isl_ast_expr;
> struct isl_ast_build;
> +struct isl_union_map;
> struct isl_pw_multi_aff;
>
> namespace polly {
> class Scop;
> class IslAst;
> +class MemoryAccess;
>
> class IslAstInfo : public ScopPass {
> public:
> - /// @brief Payload information used to annoate an ast node.
> + using MemoryAccessSet = SmallPtrSet<MemoryAccess *, 4>;
> +
> + /// @brief Payload information used to annotate an AST node.
> struct IslAstUserPayload {
> /// @brief Construct and initialize the payload.
> IslAstUserPayload()
> @@ -63,6 +67,9 @@ public:
>
> /// @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:
> @@ -82,7 +89,7 @@ public:
> /// @brief Return a copy of the AST root node.
> __isl_give isl_ast_node *getAst() const;
>
> - /// @brief Get the run conditon.
> + /// @brief Get the run condition.
> ///
> /// Only if the run condition evaluates at run-time to a non-zero value, the
> /// assumptions that have been taken hold. If the run condition evaluates to
> @@ -112,6 +119,9 @@ public:
> /// @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 broken reductions or a nullptr if not available.
> + static MemoryAccessSet *getBrokenReductions(__isl_keep isl_ast_node *Node);
> +
> ///}
>
> virtual void getAnalysisUsage(AnalysisUsage &AU) const;
> diff --git a/include/polly/Dependences.h b/include/polly/Dependences.h
> index 56f864d..3eb8362 100755
> --- a/include/polly/Dependences.h
> +++ b/include/polly/Dependences.h
> @@ -40,6 +40,7 @@ namespace polly {
>
> class Scop;
> class ScopStmt;
> +class MemoryAccess;
>
> class Dependences : public ScopPass {
> public:
> @@ -105,6 +106,16 @@ public:
> /// @brief Report if valid dependences are available.
> bool hasValidDependences();
>
> + /// @brief Return the reduction dependences caused by @p MA.
> + ///
> + /// @return The reduction dependences caused by @p MA or nullptr if None.
> + __isl_give isl_map *getReductionDependences(MemoryAccess *MA);
> +
> + /// @brief Return the reduction dependences mapped by the causing @p MA.
> + const DenseMap<MemoryAccess *, isl_map *> &getReductionDependences() const {
> + return ReductionDependences;
> + }
> +
> bool runOnScop(Scop &S);
> void printScop(raw_ostream &OS) const;
> virtual void releaseMemory();
> @@ -122,6 +133,9 @@ private:
> /// @brief The (reverse) transitive closure of reduction dependences
> isl_union_map *TC_RED = nullptr;
>
> + /// @brief Map from memory accesses to their reduction dependences.
> + DenseMap<MemoryAccess *, isl_map *> ReductionDependences;
> +
> /// @brief Collect information about the SCoP.
> void collectInfo(Scop &S, isl_union_map **Read, isl_union_map **Write,
> isl_union_map **MayWrite, isl_union_map **AccessSchedule,
> @@ -132,6 +146,10 @@ private:
>
> /// @brief Calculate the dependences for a certain SCoP.
> void calculateDependences(Scop &S);
> +
> + /// @brief Set the reduction dependences for @p MA to @p Deps.
> + void setReductionDependences(MemoryAccess *MA, __isl_take isl_map *Deps);
> +
> };
>
> } // End polly namespace.
> diff --git a/lib/Analysis/Dependences.cpp b/lib/Analysis/Dependences.cpp
> index 715b88b..1634d66 100644
> --- a/lib/Analysis/Dependences.cpp
> +++ b/lib/Analysis/Dependences.cpp
> @@ -355,6 +355,30 @@ void Dependences::calculateDependences(Scop &S) {
> DEBUG(dbgs() << "Final Wrapped Dependences:\n"; printScop(dbgs());
> dbgs() << "\n");
Could you add a comment of what you are doing here and how it is
supposed to work? (I am sure it is obvious to you, but it would be good
to document it.)
> + isl_union_map *RED_SIN = isl_union_map_empty(isl_union_map_get_space(RAW));
> + for (ScopStmt *Stmt : S) {
> + for (MemoryAccess *MA : *Stmt) {
> + if (!MA->isReductionLike())
> + continue;
> +
> + isl_set *AccDomW = isl_map_wrap(MA->getAccessRelation());
> + isl_union_map *AccRedDepU = isl_union_map_intersect_domain(
> + isl_union_map_copy(TC_RED), isl_union_set_from_set(AccDomW));
> + if (isl_union_map_is_empty(AccRedDepU) && !isl_union_map_free(AccRedDepU))
What is the free doing in the if? Is this a clever way to avoid braces?
> + continue;
> +
> + isl_map *AccRedDep = isl_map_from_union_map(AccRedDepU);
> + RED_SIN = isl_union_map_add_map(RED_SIN, isl_map_copy(AccRedDep));
> + AccRedDep = isl_map_zip(AccRedDep);
> + AccRedDep = isl_set_unwrap(isl_map_domain(AccRedDep));
> + setReductionDependences(MA, AccRedDep);
> + }
> + }
> + assert(isl_union_map_is_equal(RED_SIN, TC_RED) &&
> + "Intersecting the reduction dependence domain with the wrapped access "
> + "relation is not enough, we need to loosen the access relation also");
> + isl_union_map_free(RED_SIN);
> +
> RAW = isl_union_map_zip(RAW);
> WAW = isl_union_map_zip(WAW);
> WAR = isl_union_map_zip(WAR);
> @@ -506,6 +530,10 @@ void Dependences::releaseMemory() {
> isl_union_map_free(TC_RED);
>
> RED = RAW = WAR = WAW = TC_RED = nullptr;
> +
> + for (auto &ReductionDeps : ReductionDependences)
> + isl_map_free(ReductionDeps.second);
> + ReductionDependences.clear();
> }
>
> isl_union_map *Dependences::getDependences(int Kinds) {
> @@ -537,6 +565,16 @@ bool Dependences::hasValidDependences() {
> return (RAW != nullptr) && (WAR != nullptr) && (WAW != nullptr);
> }
>
> +isl_map *Dependences::getReductionDependences(MemoryAccess *MA) {
> + return isl_map_copy(ReductionDependences[MA]);
> +}
> +
> +void Dependences::setReductionDependences(MemoryAccess *MA, isl_map *D) {
> + assert(ReductionDependences.count(MA) == 0 &&
> + "Reduction dependences set twice!");
> + ReductionDependences[MA] = D;
> +}
> +
> void Dependences::getAnalysisUsage(AnalysisUsage &AU) const {
> ScopPass::getAnalysisUsage(AU);
> }
> diff --git a/lib/CodeGen/IslAst.cpp b/lib/CodeGen/IslAst.cpp
> index 0910835..e4f14ec 100644
> --- a/lib/CodeGen/IslAst.cpp
> +++ b/lib/CodeGen/IslAst.cpp
> @@ -100,6 +100,47 @@ struct AstBuildUserInfo {
> isl_id *LastForNodeId;
> };
>
> +static std::string getReductionOperatorStr(MemoryAccess::ReductionType RT) {
> + switch (RT) {
> + case MemoryAccess::RT_NONE:
> + llvm_unreachable("'None-type' reduction does not have an operator");
What does this mean? Should this be:
"Memory access is not a reduction"?
> + case MemoryAccess::RT_ADD:
> + return "+";
> + case MemoryAccess::RT_MUL:
> + return "*";
> + case MemoryAccess::RT_BOR:
> + return "|";
> + case MemoryAccess::RT_BXOR:
> + return "^";
> + case MemoryAccess::RT_BAND:
> + return "&";
> + }
> + llvm_unreachable("Unknown reduction type");
> + return "";
> +}
> +
> +static __isl_give isl_printer *
> +printBrokenReductions(__isl_keep isl_ast_node *Node,
> + __isl_take isl_printer *Printer,
> + __isl_take isl_ast_print_options *PrintOptions) {
> + std::map<MemoryAccess::ReductionType, std::string> ReductionClauses;
> + for (MemoryAccess *MA : *IslAstInfo::getBrokenReductions(Node)) {
> + if (!MA->isWrite())
> + continue;
> + ReductionClauses[MA->getReductionType()] +=
> + ", " + MA->getBaseName().substr(7);
Why is here a substr(7)? That looks like magic? You at least want to
hide this magic in the MemoryAccess class.
> + }
> +
> + for (const auto &ReductionClause : ReductionClauses) {
> + Printer = isl_printer_print_str(Printer, " reduction (");
> + std::string Str = getReductionOperatorStr(ReductionClause.first);
> + Str += " : " + ReductionClause.second.substr(2);
Why is here a substr(2)?
> + Printer = isl_printer_print_str(Printer, Str.c_str());
> + Printer = isl_printer_print_str(Printer, ")");
> + }
> + return Printer;
> +}
> +
> // Print a loop annotated with OpenMP or vector pragmas.
> static __isl_give isl_printer *
> printParallelFor(__isl_keep isl_ast_node *Node, __isl_take isl_printer *Printer,
> @@ -110,14 +151,14 @@ printParallelFor(__isl_keep isl_ast_node *Node, __isl_take isl_printer *Printer,
> Printer = isl_printer_start_line(Printer);
> Printer = isl_printer_print_str(Printer, "#pragma simd");
> if (Info->IsReductionParallel)
> - Printer = isl_printer_print_str(Printer, " reduction");
> + Printer = printBrokenReductions(Node, Printer, PrintOptions);
> Printer = isl_printer_end_line(Printer);
> }
> if (Info->IsOutermostParallel) {
> Printer = isl_printer_start_line(Printer);
> Printer = isl_printer_print_str(Printer, "#pragma omp parallel for");
> if (Info->IsReductionParallel)
> - Printer = isl_printer_print_str(Printer, " reduction");
> + Printer = printBrokenReductions(Node, Printer, PrintOptions);
> Printer = isl_printer_end_line(Printer);
> }
> }
> @@ -151,7 +192,7 @@ printFor(__isl_take isl_printer *Printer,
> /// (or non-zero) dependence distance on the dimension in question.
> static bool astScheduleDimIsParallel(__isl_keep isl_ast_build *Build,
> Dependences *D,
> - bool &IsReductionParallel) {
> + IslAstUserPayload *NodeInfo) {
> if (!D->hasValidDependences())
> return false;
>
> @@ -163,7 +204,20 @@ static bool astScheduleDimIsParallel(__isl_keep isl_ast_build *Build,
>
> isl_union_map *RedDeps = D->getDependences(Dependences::TYPE_TC_RED);
> if (!D->isParallel(Schedule, RedDeps))
> - IsReductionParallel = true;
> + 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;
> @@ -176,8 +230,7 @@ static void markOpenmpParallel(__isl_keep isl_ast_build *Build,
> if (BuildInfo->InParallelFor)
> return;
>
> - if (astScheduleDimIsParallel(Build, BuildInfo->Deps,
> - NodeInfo->IsReductionParallel)) {
> + if (astScheduleDimIsParallel(Build, BuildInfo->Deps, NodeInfo)) {
> BuildInfo->InParallelFor = 1;
> NodeInfo->IsOutermostParallel = 1;
> }
> @@ -225,8 +278,7 @@ astBuildAfterFor(__isl_take isl_ast_node *Node, __isl_keep isl_ast_build *Build,
> if (Info->IsOutermostParallel)
> BuildInfo->InParallelFor = 0;
> if (IsInnermost)
> - if (astScheduleDimIsParallel(Build, BuildInfo->Deps,
> - Info->IsReductionParallel))
> + if (astScheduleDimIsParallel(Build, BuildInfo->Deps, Info))
> Info->IsInnermostParallel = 1;
> if (!Info->Build)
> Info->Build = isl_ast_build_copy(Build);
> @@ -383,6 +435,12 @@ isl_union_map *IslAstInfo::getSchedule(__isl_keep isl_ast_node *Node) {
> return Payload ? isl_ast_build_get_schedule(Payload->Build) : nullptr;
> }
>
> +IslAstInfo::MemoryAccessSet *
> +IslAstInfo::getBrokenReductions(__isl_keep isl_ast_node *Node) {
> + IslAstUserPayload *Payload = getNodePayload(Node);
> + return Payload ? &Payload->BrokenReductions : nullptr;
> +}
> +
> void IslAstInfo::printScop(raw_ostream &OS) const {
> isl_ast_print_options *Options;
> isl_ast_node *RootNode = getAst();
> diff --git a/test/Isl/Ast/reduction_clauses_multidimensional_access.ll b/test/Isl/Ast/reduction_clauses_multidimensional_access.ll
> new file mode 100644
> index 0000000..dba85cd
> --- /dev/null
> +++ b/test/Isl/Ast/reduction_clauses_multidimensional_access.ll
> @@ -0,0 +1,68 @@
> +; RUN: opt %loadPolly -polly-delinearize -polly-ast -polly-ast-detect-parallel -analyze < %s | FileCheck %s
> +;
> +; CHECK: #pragma omp parallel for reduction (^ : sum)
> +; void f(int N, int M, int P, int sum[P][M]) {
> +; for (int i = 0; i < N; i++)
> +; for (int j = 0; j < P; j++)
> +; CHECK: #pragma simd
> +; for (int k = 0; k < M; k++)
> +; sum[j][k] ^= j;
> +; }
Very nice test cases!
Tobias
More information about the llvm-commits
mailing list