sdfg::tiles::GPUTileTarget Class Reference

The built-in GPU backend for the tile algebra, shared by CUDA and ROCm. More...

#include <sdfg/targets/gpu/gpu_tile_target.h>

Inheritance diagram for sdfg::tiles::GPUTileTarget:

Public Member Functions
	GPUTileTarget (unsigned lane_width, std::string offload_value, data_flow::ImplementationType implementation_type, size_t min_async_bytes=4)
std::optional< 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).
Space	space (Level level) const override
	The GPU scratchpad hierarchy: Device->Global, Group->Shared, Subgroup->Register.
bool	supports_cooperative_staging (const structured_control_flow::ScheduleType &sched) const override
	Whether sched can host a block-cooperative shared-memory staging copy driven by its own parallel threads (i.e.
types::StorageType	storage_type (Space space) const override
	The concrete buffer that realizes an abstract cooperation Space on this target: e.g.
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).
std::optional< tiles::BankModel >	bank_model (tiles::Space space) const override
	Shared memory (NVIDIA and CDNA LDS): 32 banks of 4 bytes.
data_flow::ImplementationType	implementation_type () const override
	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.
size_t	min_async_bytes () const override
	The narrowest per-lane async (CpAsync) transfer in bytes; narrower copies stay synchronous.
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 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	max_tile_elements () const
	The most elements a staging buffer local to one of this target's loops may hold.

Detailed Description

The built-in GPU backend for the tile algebra, shared by CUDA and ROCm.

Maps an offload schedule's target level to a cooperation Level and the abstract Space tiers to NVIDIA/CDNA buffers. lane_width is the warp/wavefront size (32 on NVIDIA, 64 on CDNA); offload_value is the *_Offload schedule value this instance owns (the legacy fused schedule is the other value it is registered under).

Constructor & Destructor Documentation

GPUTileTarget()

sdfg::tiles::GPUTileTarget::GPUTileTarget ( unsigned lane_width, std::string offload_value, data_flow::ImplementationType implementation_type, size_t min_async_bytes = 4 )

inline

Member Function Documentation

bank_model()

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

inlineoverridevirtual

Shared memory (NVIDIA and CDNA LDS): 32 banks of 4 bytes.

Reimplemented from sdfg::tiles::TileTarget.

classify()

std::optional< AxisSchedule > sdfg::tiles::GPUTileTarget::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::tiles::GPUTileTarget::implementation_type ( ) const

inlineoverridevirtual

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.

Implements sdfg::tiles::TileTarget.

lane_width()

unsigned sdfg::tiles::GPUTileTarget::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.

min_async_bytes()

size_t sdfg::tiles::GPUTileTarget::min_async_bytes ( ) const

inlineoverridevirtual

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

Reimplemented from sdfg::tiles::TileTarget.

space()

Space sdfg::tiles::GPUTileTarget::space ( Level level ) const

inlineoverridevirtual

The GPU scratchpad hierarchy: Device->Global, Group->Shared, Subgroup->Register.

Implements sdfg::tiles::TileTarget.

storage_type()

types::StorageType sdfg::tiles::GPUTileTarget::storage_type ( 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::tiles::GPUTileTarget::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/gpu/gpu_tile_target.h