sdfg::omp::OMPTileTarget Class Reference

#include <sdfg/targets/omp/plugin.h>

Inheritance diagram for sdfg::omp::OMPTileTarget:

Public Member Functions
std::optional< tiles::AxisSchedule >	classify (const structured_control_flow::ScheduleType &sched) const override
	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).
types::StorageType	storage_type (tiles::Space space) const override
	The concrete buffer that realizes an abstract cooperation Space on this target: e.g.
tiles::Space	space (tiles::Level) const override
	A host thread pool has no on-chip scratchpad: every level cooperates through global memory.
bool	supports_cooperative_staging (const structured_control_flow::ScheduleType &) const override
	Whether sched can host a block-cooperative shared-memory staging copy driven by its own parallel threads (i.e.
unsigned	lane_width () const override
	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).
data_flow::ImplementationType	implementation_type () const override
	A host thread pool has no cp.async / vector-copy / pipeline library nodes.
Public Member Functions inherited from sdfg::tiles::TileTarget
virtual	~TileTarget ()=default
bool	has_scratchpad () const
	Whether this target exposes any on-chip cooperative tier below global memory (a scratchpad hierarchy), derived from space.
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.

Member Function Documentation

classify()

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

inlineoverridevirtual

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.

Implements sdfg::tiles::TileTarget.

implementation_type()

data_flow::ImplementationType sdfg::omp::OMPTileTarget::implementation_type ( ) const

inlineoverridevirtual

A host thread pool has no cp.async / vector-copy / pipeline library nodes.

Implements sdfg::tiles::TileTarget.

lane_width()

unsigned sdfg::omp::OMPTileTarget::lane_width ( ) const

inlineoverridevirtual

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.

Implements sdfg::tiles::TileTarget.

space()

tiles::Space sdfg::omp::OMPTileTarget::space ( tiles::Level ) const

inlineoverridevirtual

A host thread pool has no on-chip scratchpad: every level cooperates through global memory.

Implements sdfg::tiles::TileTarget.

storage_type()

types::StorageType sdfg::omp::OMPTileTarget::storage_type ( tiles::Space space ) const

inlineoverridevirtual

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.

Implements sdfg::tiles::TileTarget.

supports_cooperative_staging()

bool sdfg::omp::OMPTileTarget::supports_cooperative_staging ( const structured_control_flow::ScheduleType & sched ) const

inlineoverridevirtual

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.

Implements sdfg::tiles::TileTarget.

---

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

• opt/include/sdfg/targets/omp/plugin.h