|
docc 0.8.3 (0d44999b)
C++ API Reference
|
Software-pipeline a sequential panel loop that cooperatively stages a shared-memory tile each iteration (the structure LocalStorage emits). More...
#include <sdfg/tiles/transformations/software_pipelining.h>
Public Member Functions | |
| SoftwarePipelining (structured_control_flow::StructuredLoop &loop, size_t stages=2, bool single_operand=false, bool register_staged=false) | |
| virtual std::string | name () const override |
| Get the name of this transformation. | |
| virtual bool | can_be_applied (builder::StructuredSDFGBuilder &builder, analysis::AnalysisManager &analysis_manager) override |
| Check if this transformation can be safely applied. | |
| virtual void | apply (builder::StructuredSDFGBuilder &builder, analysis::AnalysisManager &analysis_manager) override |
| Apply this transformation to the SDFG. | |
| virtual void | to_json (nlohmann::json &j) const override |
| Serialize this transformation to JSON. | |
| Public Member Functions inherited from sdfg::transformations::Transformation | |
| virtual | ~Transformation ()=default |
| virtual void | set_report (PassReportConsumer *report) |
| Set the report consumer for transformation feedback. | |
| virtual bool | try_apply (builder::StructuredSDFGBuilder &builder, analysis::AnalysisManager &analysis_manager) |
| Attempt to apply the transformation if possible. | |
Static Public Member Functions | |
| static std::set< std::string > | copied_in_buffers (const Function &function, structured_control_flow::ControlFlowNode &root) |
The async-stageable buffers the copy-in nodes under root fill: the containers a pipeline would multi-buffer. | |
| static std::vector< structured_control_flow::StructuredLoop * > | staging_panels (const Function &function, structured_control_flow::ControlFlowNode &root) |
The candidate panel loops under root this transformation can target. | |
| static SoftwarePipelining | from_json (builder::StructuredSDFGBuilder &builder, const nlohmann::json &j) |
Additional Inherited Members | |
| Protected Attributes inherited from sdfg::transformations::Transformation | |
| PassReportConsumer * | report_ = nullptr |
Software-pipeline a sequential panel loop that cooperatively stages a shared-memory tile each iteration (the structure LocalStorage emits).
Targets a sequential loop whose body is [ barrier; cooperative-copy(global -> shared); barrier; compute(reads shared) ] and rewrites it into an stages -deep pipeline so the next panel's global load (issued via cp.async) overlaps the current panel's compute:
prologue: async-copy panel 0 -> buf[0]; commit for p in [0, n): if p + stages-1 < n: async-copy panel p+stages-1 -> buf[(p+stages-1)stages]; commit wait_prior(stages-1); barrier compute from buf[p % stages]
The staged buffer gains a leading [stages] axis indexed by p % stages. The staging copy's tiles::TileCopyNode has its atom flipped to cp.async (its plan offset biased per stage) and gains a PipelineCommitNode; a PipelineWaitNode fences each panel's reads. Only fires inside a map whose tile target supports asynchronous staging of the copied-in buffers (tiles::TileTarget::supports_async_staging; on GPUs the shared tier), on a compile-time-constant panel count >= stages; buffers written by the compute stay single-buffered.
|
explicit |
|
overridevirtual |
Apply this transformation to the SDFG.
| builder | The SDFG builder |
| analysis_manager | The analysis manager (may be invalidated) |
Implements sdfg::transformations::Transformation.
|
overridevirtual |
Check if this transformation can be safely applied.
| builder | The SDFG builder |
| analysis_manager | The analysis manager providing necessary analyses |
Implements sdfg::transformations::Transformation.
|
static |
The async-stageable buffers the copy-in nodes under root fill: the containers a pipeline would multi-buffer.
Empty when root stages nothing (so it is not a pipelineable panel).
A buffer qualifies when a cooperative tiles::TileCopyNode copies into it (its _dst pointer) and the enclosing map's tile target can asynchronously multi-buffer its storage (tiles::TileTarget::supports_async_staging; on GPUs the shared tier). This is the single source of truth for "which buffers get staged", shared by can_be_applied, apply, and search spaces.
|
static |
|
overridevirtual |
Get the name of this transformation.
Implements sdfg::transformations::Transformation.
|
static |
The candidate panel loops under root this transformation can target.
A panel is the nearest loop enclosing a cooperative copy-in (a tiles::TileCopyNode that LocalStorage emits) that fills a buffer the loop's enclosing tile target can asynchronously multi-buffer (tiles::TileTarget::supports_async_staging; on GPUs the shared tier). Only sequential loops qualify: a copy-in sitting directly in a parallel Map has no cross-iteration order to pipeline, so it nominates no panel. Returned in first-appearance order. This is the single source of truth for "what can be pipelined", so a search space can enumerate candidates without re-deriving the scan.
|
overridevirtual |
Serialize this transformation to JSON.
| j | JSON object to populate with transformation description |
| std::logic_error | if not implemented |
Implements sdfg::transformations::Transformation.