sdfg::tiles::TileTarget Class Reference

#include <sdfg/tiles/tile_target.h>

Inheritance diagram for sdfg::tiles::TileTarget:

Public Member Functions
virtual	~TileTarget ()=default
virtual std::optional< AxisSchedule >	classify (const structured_control_flow::ScheduleType &sched) const =0
	Classify one loop's sched into target-neutral cooperation facts (AxisSchedule: level, spatial axis, parallel size, sync), or std::nullopt when the schedule does not shape storage (a sequential loop).
virtual Space	space (Level level) const =0
	The memory Space backing cooperation at each Level on this target: a GPU maps Group -> Shared, Subgroup -> Register; a flat CPU maps every level to Global; an exotic scratchpad target maps whichever levels it materializes on-chip.
bool	has_scratchpad () const
	Whether this target exposes any on-chip cooperative tier below global memory (a scratchpad hierarchy), derived from space.
virtual bool	supports_cooperative_staging (const structured_control_flow::ScheduleType &sched) const =0
	Whether sched can host a block-cooperative shared-memory staging copy driven by its own parallel threads (i.e.
virtual types::StorageType	storage_type (Space space) const =0
	The concrete buffer that realizes an abstract cooperation Space on this target: e.g.
virtual unsigned	lane_width () const =0
	The SIMD lane / subgroup width, used only to keep a subgroup's cooperative stores bank-conflict-free (32 on NVIDIA, 64 on CDNA, 1 where there is no SIMD cooperation).
virtual std::optional< BankModel >	bank_model (Space space) const
	The banks of the memory realizing space, if it is banked (GPU shared memory).
virtual bool	supports_async_staging (const types::StorageType &storage) const
	Whether staging buffers of storage can be multi-buffered so that a loop prefetches its next tile asynchronously while computing on the current one (SoftwarePipelining).
virtual size_t	min_async_bytes () const
	The narrowest per-lane async (CpAsync) transfer in bytes; narrower copies stay synchronous.
virtual size_t	max_tile_elements () const
	The most elements a staging buffer local to one of this target's loops may hold.
virtual data_flow::ImplementationType	implementation_type () const =0
	The implementation selector this target stamps onto its tile codegen library nodes (cp.async / vector copy / pipeline commit+wait) so the owning target's dispatcher is chosen at codegen.

Constructor & Destructor Documentation

~TileTarget()

virtual sdfg::tiles::TileTarget::~TileTarget ( )

virtualdefault

Member Function Documentation

bank_model()

virtual std::optional< BankModel > sdfg::tiles::TileTarget::bank_model ( Space space ) const

inlinevirtual

The banks of the memory realizing space, if it is banked (GPU shared memory).

Drives the bank-conflict policy of buffers staged there; unbanked memory returns nullopt.

Reimplemented in sdfg::tiles::GPUTileTarget.

classify()

virtual std::optional< AxisSchedule > sdfg::tiles::TileTarget::classify ( const structured_control_flow::ScheduleType & sched ) const

pure virtual

Classify one loop's sched into target-neutral cooperation facts (AxisSchedule: level, spatial axis, parallel size, sync), or std::nullopt when the schedule does not shape storage (a sequential loop).

Replaces the target-string switch inside AxisSchedule::classify.

Implemented in sdfg::omp::OMPTileTarget, and sdfg::tiles::GPUTileTarget.

has_scratchpad()

bool sdfg::tiles::TileTarget::has_scratchpad ( ) const

inline

Whether this target exposes any on-chip cooperative tier below global memory (a scratchpad hierarchy), derived from space.

classify stamps it onto every axis so the tile algebra can tell a device-wide axis that sits atop a scratchpad (GPU grid) from a flat host axis, without naming any target.

implementation_type()

virtual data_flow::ImplementationType sdfg::tiles::TileTarget::implementation_type ( ) const

pure virtual

The implementation selector this target stamps onto its tile codegen library nodes (cp.async / vector copy / pipeline commit+wait) so the owning target's dispatcher is chosen at codegen.

CUDA returns "CUDA", ROCm "ROCM"; a target without these nodes returns data_flow::ImplementationType_NONE. This is the seam the pipelining/vectorization transformations reach through instead of hard-coding a backend.

Implemented in sdfg::omp::OMPTileTarget, and sdfg::tiles::GPUTileTarget.

lane_width()

virtual unsigned sdfg::tiles::TileTarget::lane_width ( ) const

pure virtual

The SIMD lane / subgroup width, used only to keep a subgroup's cooperative stores bank-conflict-free (32 on NVIDIA, 64 on CDNA, 1 where there is no SIMD cooperation).

Replaces gpu::gpu_warp_size in the padding heuristic.

Implemented in sdfg::omp::OMPTileTarget, and sdfg::tiles::GPUTileTarget.

max_tile_elements()

virtual size_t sdfg::tiles::TileTarget::max_tile_elements ( ) const

inlinevirtual

The most elements a staging buffer local to one of this target's loops may hold.

min_async_bytes()

virtual size_t sdfg::tiles::TileTarget::min_async_bytes ( ) const

inlinevirtual

The narrowest per-lane async (CpAsync) transfer in bytes; narrower copies stay synchronous.

Reimplemented in sdfg::tiles::GPUTileTarget.

space()

virtual Space sdfg::tiles::TileTarget::space ( Level level ) const

pure virtual

The memory Space backing cooperation at each Level on this target: a GPU maps Group -> Shared, Subgroup -> Register; a flat CPU maps every level to Global; an exotic scratchpad target maps whichever levels it materializes on-chip.

classify stamps each axis with space(its level). This is the single per-level, per-target memory fact the tile algebra reads.

Implemented in sdfg::omp::OMPTileTarget, and sdfg::tiles::GPUTileTarget.

storage_type()

virtual types::StorageType sdfg::tiles::TileTarget::storage_type ( Space space ) const

pure virtual

The concrete buffer that realizes an abstract cooperation Space on this target: e.g.

Shared -> NV_Shared / TT_L1, Global -> NV_Global, Register -> a thread-private stack/register block. Replaces the Space-to-StorageType switch in LocalStorage.

Implemented in sdfg::omp::OMPTileTarget, and sdfg::tiles::GPUTileTarget.

supports_async_staging()

virtual bool sdfg::tiles::TileTarget::supports_async_staging ( const types::StorageType & storage ) const

inlinevirtual

Whether staging buffers of storage can be multi-buffered so that a loop prefetches its next tile asynchronously while computing on the current one (SoftwarePipelining).

Defaults to the on-chip shared tier of a scratchpad target.

supports_cooperative_staging()

virtual bool sdfg::tiles::TileTarget::supports_cooperative_staging ( const structured_control_flow::ScheduleType & sched ) const

pure virtual

Whether sched can host a block-cooperative shared-memory staging copy driven by its own parallel threads (i.e.

a genuine offload schedule), as opposed to a legacy fused whole-kernel schedule that cooperates at group level but cannot carry a separate copy Map. Distinguishes the two where classify alone maps both to Level::Group.

Implemented in sdfg::omp::OMPTileTarget, and sdfg::tiles::GPUTileTarget.

---

The documentation for this class was generated from the following file:

• opt/include/sdfg/tiles/tile_target.h