Local-storage transformation.
More...
#include <sdfg/tiles/transformations/local_storage.h>
|
| | LocalStorage (structured_control_flow::StructuredLoop &loop, const data_flow::AccessNode &access_node, bool swizzle_layout=false, bool lane_contiguous=false, bool transpose_layout=false) |
| | Construct a local-storage transformation for access_node's container within loop.
|
| std::string | name () const override |
| | Get the name of this transformation.
|
| bool | can_be_applied (builder::StructuredSDFGBuilder &builder, analysis::AnalysisManager &analysis_manager) override |
| | Precondition check: verifies the container is a pointer used in the loop, resolves a single tile group with integer extents, derives the storage space, and populates tile_info_ — rejecting any schedule the apply path cannot yet handle.
|
| void | apply (builder::StructuredSDFGBuilder &builder, analysis::AnalysisManager &analysis_manager) override |
| | Allocate the buffer, emit the derived copies, and redirect loop-body accesses to the buffer.
|
| void | to_json (nlohmann::json &j) const override |
| | JSON serialization.
|
| const std::string & | local_container () const |
| | Name of the created local buffer (valid after apply()).
|
| const TileInfo & | tile_info () const |
| | Tile info (valid after can_be_applied() returns true).
|
| const types::StorageType & | storage_type () const |
| | Storage type of the local buffer.
|
| 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.
|
Local-storage transformation.
Relocates a compile-time-bounded region (a tile) of a large pointer container into a small, densely packed, contiguous local buffer whose storage space is derived from the schedule (see tiles::LocalityPlan::required_space).
◆ LocalStorage()
Construct a local-storage transformation for access_node's container within loop.
Direction and storage space are derived by can_be_applied().
- Parameters
-
| swizzle_layout | A bank-conflict-free NV_Shared tile uses an XOR swizzle of the inner index instead of padding — no wasted columns and the layout ldmatrix needs. Requires a power-of-two inner block; falls back to padding otherwise. |
| lane_contiguous | The NV_Shared tile is laid out flat and thread-linear (slots folded, no padding), staged by a full-block cooperative copy. Required by the CDNA async global->LDS DMA (global_load_lds). |
| transpose_layout | Store a cooperative (no-slot) NV_Shared tile column-major (its tile axes reversed), so consumers read it transposed without a separate pass. A pure affine relabelling of storage. |
◆ apply()
◆ can_be_applied()
Precondition check: verifies the container is a pointer used in the loop, resolves a single tile group with integer extents, derives the storage space, and populates tile_info_ — rejecting any schedule the apply path cannot yet handle.
Implements sdfg::transformations::Transformation.
◆ from_json()
◆ has_side_effect()
True if any library node in loop's body has side effects (may touch memory outside the tracked memlets, so localization must bail).
◆ local_container()
| const std::string & sdfg::transformations::LocalStorage::local_container |
( |
| ) |
const |
|
inline |
Name of the created local buffer (valid after apply()).
◆ name()
| std::string sdfg::transformations::LocalStorage::name |
( |
| ) |
const |
|
inlineoverridevirtual |
◆ storage_type()
Storage type of the local buffer.
◆ tile_info()
| const TileInfo & sdfg::transformations::LocalStorage::tile_info |
( |
| ) |
const |
|
inline |
◆ to_json()
| void sdfg::transformations::LocalStorage::to_json |
( |
nlohmann::json & | j | ) |
const |
|
overridevirtual |
The documentation for this class was generated from the following file: