Classes
class	GPUTileTarget
	The built-in GPU backend for the tile algebra, shared by CUDA and ROCm. More...
struct	ReductionScheduleProposal
	A schedule-only change; loop domains, access expressions, and nesting stay fixed. More...
struct	ReductionLoopDomain
	One projected affine loop domain; geometry queries consume domains deepest first. More...
struct	ReductionBufferInfo
	One accumulator's logical footprint and allocation contribution. More...
struct	ReductionKernelInfo
	Aggregate reduction-owned shared storage; excludes unrelated staging buffers. More...
class	ReductionBufferAnalysis
	Infer reduction footprints and verify explicitly materialized buffers. More...
class	TileAnalysis
	Schedule-aware tile analysis: for a (loop scope, container) that MemoryLayoutAnalysis resolves to a bounded tile, produces a tiles::Tile (geometry, thread/value partition, copy direction) for memory-level transformations to consume instead of re-deriving. More...
struct	BankModel
	A banked on-chip memory: banks banks of bank_bytes each, every bank serving one word per cycle. More...
struct	PackedBuffer
	The packed local buffer as a value: its nested-array shape (axes), the matching per-axis element address (subset), and the scalar-offset layout. More...
class	PipelineCommitNode
class	PipelineWaitNode
class	PipelineCommitNodeSerializer
class	PipelineWaitNodeSerializer
class	TileCopyNode
	A tiled copy carried as a value TiledCopy plan. More...
class	TileCopyNodeSerializer
class	TileCopyNodeDispatcher
	Reference (sequential, scalar) dispatcher for TileCopyNode. More...
class	LocalityPlan
	How a tile relates to its enclosing parallel loop nest, and the storage space that relationship demands. More...
struct	Swizzle
	XOR-swizzle as a composition functor over a Layout. More...
struct	ComposedLayout
	A layout with a swizzle on its output offset (S o L) — the bank-conflict ConflictPolicy::Swizzle. More...
class	AxisSchedule
	Target-neutral cooperation facts derived from one loop's schedule. More...
class	TileAxis
	One enclosing parallel loop of a tile: its schedule facts plus the tile-role (Private when its indvar addresses the tile base, Cooperative otherwise). More...
class	Tile
	A container's staged region: source maps a tile coordinate to the global element, axes is the thread/value partition, and the read/write flags give the copy direction. More...
class	TileTarget
class	TileTargetRegistry
	Maps a schedule value() to the TileTarget that owns it. More...
struct	TiledCopy
struct	TileGuard
	The ragged-tile boundary predicate, stored per-dim. More...
class	TileGuardNormalization
	Re-discharge every TileCopyNode's ragged boundary guard. More...

Typedefs
using	Layout = math::tensor::TensorLayout
	The tiles geometry type is math::tensor::TensorLayout — the single source of truth for a (shape, strides, offset) affine element map.
using	GuardedLoadFn
	Renders a guarded (ragged-tile) copy-in load for the cooperative copy loop: a statement declaring type var that holds the `bytes` -wide value at (base)[index] (an element index into the flat `elem_bytes` -sized element pointer `base`), or zero where the in-scope __tc_ok predicate is false.

Enumerations
enum class	ReductionBufferStatus { Exact , ConservativeBound , Unsupported }
	How precisely the analysis can describe a reduction's storage requirements. More...
enum class	Placement { RowMajor , ColMajor , Flat }
enum class	ConflictPolicy { None , PadRow , PadSlot , Swizzle }
enum class	Level { Device , Group , Subgroup }
	Cooperation scope of a parallel axis, coarsest to finest — target-neutral tiers (OpenCL/SYCL vocabulary). More...
enum class	Space { Global , Shared , Register }
	Memory that backs cooperation at a level: Device/CPU -> Global (atomics/heap), Group -> Shared, Subgroup/none -> Register. More...
enum class	Role { Private , Cooperative }
	Whether the tile is per-iteration-private or shared across a parallel axis. More...
enum class	CopyAtom { ScalarSync , VectorSync , CpAsync , CoalesceSync }
	The per-lane transfer primitive a tiled copy lowers to. More...
enum class	CopyDirection { In , Out }
	Which way the tile moves between the flat (global) side and the packed buffer. More...
enum class	CopyPhase { Full , LoadRegs , StoreRegs }
	Which half of a register-staged copy a node performs. More...

Functions
bool	is_reduction_accumulator (structured_control_flow::StructuredLoop &loop, const std::string &container, analysis::AnalysisManager &analysis_manager)
	Whether `container` is the accumulator of a Reduce enclosing, nested within, or equal to `loop`.
bool	collect_reduction_owners (structured_control_flow::StructuredLoop &loop, const std::string &container, analysis::AnalysisManager &analysis_manager, std::vector< structured_control_flow::Reduce * > &out)
	Collect the non-cooperative Reduce nodes (`loop` or a descendant) that reduce into `container`, for accumulator privatization.
std::vector< structured_control_flow::Reduce * >	collect_grid_reduction_owners (structured_control_flow::StructuredLoop &loop, const std::string &container, analysis::AnalysisManager &analysis_manager)
	Collect grid-parallel (GPU-offloaded) ancestor Reduce nodes reducing into `container`.
bool	is_constant_bounded (const analysis::MemoryTileGroup *group)
	True iff `group` is a real tile whose extents are all compile-time integer constants (the constant-bounded localization precondition).
const std::vector< analysis::MemoryTileGroup > *	tile_groups (structured_control_flow::StructuredLoop &loop, const std::string &container, analysis::AnalysisManager &analysis_manager)
	All tile groups MemoryLayoutAnalysis formed for `container` at `loop` (nullptr if none).
symbolic::Expression	tile_element_count (const analysis::MemoryTileGroup *group)
	Scalar slots a packed buffer for `group` would occupy (product of extents): a compile-time integer iff is_constant_bounded, null if any extent is unbounded, 0 for a null group.
const analysis::MemoryTileGroup *	localizable_tile (structured_control_flow::StructuredLoop &loop, const std::string &container, analysis::AnalysisManager &analysis_manager)
	The single localizable tile group of `container` at `loop`, or nullptr.
size_t	access_cycles (const BankModel &model, size_t elem_bytes, const std::vector< int64_t > &strides, const data_flow::SubgroupAccess &access)
	Cycles the banked memory spends serving `access` when the operand's elements (of `elem_bytes`) sit at element offset sum(coord[d] * strides[d]).
size_t	slot_pad (const BankModel &model, size_t block, size_t coop_span)
	ConflictPolicy::PadSlot padding (elements) of a per-slot block of `block` elements whose slots are accessed by `coop_span` consecutive lanes of a subgroup: makes the slot stride congruent to `coop_span` modulo the bank count, rounded up to a multiple of 4 elements so blocks stay vector-addressable.
size_t	row_pad (const BankModel &model, size_t elem_bytes, PackedBuffer buffer, size_t align_bytes, const std::function< std::optional< size_t >(const PackedBuffer &)> &cost)
	The ConflictPolicy::PadRow padding (elements) of `buffer` minimizing `cost`, the access cycles of a candidate (nullopt rejects it).
bool	provably_multiple_of (symbolic::Expression e, long long m, const TileCopyNode &node)
	True if `e` is a multiple of `m` for all values of its symbols: the indvar of each loop enclosing `node` is rewritten as init + stride * t; any other symbol is an unknown integer.
bool	guard_axis_aligned (const TileCopyNode &node, size_t axis, long long factor)
	Every guard dim on tile `axis` has a base and exclusive bound that are multiples of `factor`, so a factor-aligned run of coordinates along it is entirely in or entirely out of bounds.
size_t	widest_legal_transfer (const TileCopyNode &node, CopyAtom atom, const std::vector< size_t > &widths)
	The first of `widths` (bytes, widest first) at which each lane of `node` can move one contiguous multi-element run under `atom`, or 0.
std::string	emit_tile_copy_stmt (codegen::LanguageExtension &language_extension, const TileCopyNode &node, const std::string &dst_expr, const std::string &src_expr, types::PrimitiveType elem, const symbolic::Expression &index)
	The single-element copy statement dst[dst_geo(index)] = src[src_geo(index)]; for the tile coordinate `index`, addressing both bare-pointer sides as flat `elem` pointers via the plan's layouts.
std::string	emit_predicated_guarded_load (const std::string &type, const std::string &var, const std::string &base, const std::string &index)
	Vendor-agnostic GuardedLoadFn: a predicated load into a zero-initialized value (CUDA: @p LDG; AMDGPU: an exec-masked load).
void	emit_cooperative_copy_loop (codegen::LanguageExtension &language_extension, codegen::PrettyPrinter &stream, const TileCopyNode &node, const std::string &dst_expr, const std::string &src_expr, types::PrimitiveType elem, const std::function< std::string(const std::string &dst_addr, const std::string &src_addr, size_t bytes)> &async_stmt={}, const GuardedLoadFn &guarded_load={})
	Emit the block-cooperative staging loop for `node` into `stream:` every thread of the enclosing kernel block strides the flat tile from its flat index by the block thread count.
structured_control_flow::StructuredLoop *	find_block_scheduled_descendant (structured_control_flow::StructuredLoop &loop, analysis::AnalysisManager &analysis_manager)
	First block-level GPU-offloaded loop in `loop's` body (a cooperative copy / consumer axis for enclosing-scope staging), or nullptr.
void	register_tiles_plugin (plugins::Context &context)
void	register_tiles_plugin ()
symbolic::MultiExpression	delinearize_rowmajor (const symbolic::Expression &flat, const symbolic::MultiExpression &sizes)
	Row-major delinearization of a flat index into per-dim coordinates (dim 0 slowest), matching the buffer's tile linearization so both plan sides agree.
symbolic::Expression	simplify_tile_bound (const symbolic::Expression &max, const symbolic::Expression &base, const symbolic::Expression &size)
	Simplify a tile's inclusive bound `max` for coordinates base + i, 0 <= i < size, preserving the predicate base + i <= max on that range: drops min/max arguments dominated by another and min arguments at or above the tile's last coordinate (they never bind).
bool	coalesce_sweep_order (TiledCopy &plan, TileGuard &guard)
	Move the source's unit-stride tile dim last (lane-fastest in the row-major sweep), permuting both plan sides and the guard axes alike.

Variables
data_flow::LibraryNodeCode	LibraryNodeType_PipelineCommit {"pipeline_commit"}
data_flow::LibraryNodeCode	LibraryNodeType_PipelineWait {"pipeline_wait"}
data_flow::LibraryNodeCode	LibraryNodeType_TileCopy {"tile_copy"}

Typedef Documentation

◆ GuardedLoadFn

using sdfg::tiles::GuardedLoadFn

Initial value:

 std::function<std::string(
    const std::string& type,
    const std::string& var,
    const std::string& base,
    const std::string& index,
    size_t elem_bytes,
    size_t bytes
)>

Renders a guarded (ragged-tile) copy-in load for the cooperative copy loop: a statement declaring type var that holds the bytes -wide value at (base)[index] (an element index into the flat elem_bytes -sized element pointer base), or zero where the in-scope __tc_ok predicate is false.

It must not select on the loaded value: such a select forces a wait right after the load, which serializes register-staged (software-pipelined) prefetches.

◆ Layout

using sdfg::tiles::Layout = math::tensor::TensorLayout

The tiles geometry type is math::tensor::TensorLayout — the single source of truth for a (shape, strides, offset) affine element map.

A tile coordinate x addresses offset + sum_k x_k * strides_k via TensorLayout::resolve_element.

Enumeration Type Documentation

◆ ConflictPolicy

enum class sdfg::tiles::ConflictPolicy

strong

Enumerator
None
PadRow	append PackedBuffer::pad elements to every stored innermost row
PadSlot	flatten each per-slot block and widen its stride by PackedBuffer::pad elements
Swizzle	flatten each per-slot block and XOR its index with the slot index

◆ CopyAtom

enum class sdfg::tiles::CopyAtom

strong

The per-lane transfer primitive a tiled copy lowers to.

Enumerator
ScalarSync	one element per lane, synchronous store
VectorSync	a contiguous 4/8/16-byte vector per lane, synchronous
CpAsync	a contiguous 4/8/16-byte async global->shared cp.async
CoalesceSync	2-D 16/32-bit copy into a column-major buffer that keeps global reads coalesced: each lane loads four 8/16-byte source rows of a 4x4 block and regroups them into four buffer rows.

◆ CopyDirection

enum class sdfg::tiles::CopyDirection

strong

Which way the tile moves between the flat (global) side and the packed buffer.

Enumerator
In	read the global tile into the local buffer (copy-in / stage)
Out	write the local buffer back to the global tile (copy-out / writeback)

◆ CopyPhase

enum class sdfg::tiles::CopyPhase

strong

Which half of a register-staged copy a node performs.

A staged pair shares one plan and per-thread register array (32-bit words, indexed by the copy step), so the load can be issued early and the buffer store deferred.

Enumerator
Full	source -> destination
LoadRegs	source -> registers (_dst is the register array)
StoreRegs	registers -> destination (_src is the register array)

◆ Level

enum class sdfg::tiles::Level

strong

Cooperation scope of a parallel axis, coarsest to finest — target-neutral tiers (OpenCL/SYCL vocabulary).

Each target maps its schedule onto them via AxisSchedule::classify: GPU grid / CPU threads -> Device, GPU block / Tenstorrent core-workers -> Group, GPU warp / wavefront -> Subgroup.

Enumerator
Device
Group
Subgroup

◆ Placement

enum class sdfg::tiles::Placement

strong

Enumerator
RowMajor	[slot ++ tile], last tile dim fastest
ColMajor	[slot ++ reversed tile], first tile dim fastest
Flat	one flat axis over [slot ++ tile] (lane-contiguous staging)

◆ ReductionBufferStatus

enum class sdfg::tiles::ReductionBufferStatus

strong

How precisely the analysis can describe a reduction's storage requirements.

Enumerator
Exact	Exact layout and allocation costs, suitable for materialization.
ConservativeBound	Allocation upper bounds only; not a materializable layout.
Unsupported	No supported footprint; consult the diagnostic.

◆ Role

enum class sdfg::tiles::Role

strong

Whether the tile is per-iteration-private or shared across a parallel axis.

Enumerator
Private
Cooperative

◆ Space

enum class sdfg::tiles::Space

strong

Memory that backs cooperation at a level: Device/CPU -> Global (atomics/heap), Group -> Shared, Subgroup/none -> Register.

Enumerator
Global
Shared
Register

Function Documentation

◆ access_cycles()

size_t sdfg::tiles::access_cycles	(	const BankModel &	model,
		size_t	elem_bytes,
		const std::vector< int64_t > &	strides,
		const data_flow::SubgroupAccess &	access )

Cycles the banked memory spends serving access when the operand's elements (of elem_bytes) sit at element offset sum(coord[d] * strides[d]).

A lane whose elements are adjacent and aligned issues one vector access of at most 16 B, otherwise one access per element. Lanes are served in phases that fill every bank once; a phase costs the largest number of distinct words any bank serves (lanes reading the same word share it).

◆ coalesce_sweep_order()

bool sdfg::tiles::coalesce_sweep_order	(	TiledCopy &	plan,
		TileGuard &	guard )

Move the source's unit-stride tile dim last (lane-fastest in the row-major sweep), permuting both plan sides and the guard axes alike.

Returns true if the order changed.

◆ collect_grid_reduction_owners()

std::vector< structured_control_flow::Reduce * > sdfg::tiles::collect_grid_reduction_owners	(	structured_control_flow::StructuredLoop &	loop,
		const std::string &	container,
		analysis::AnalysisManager &	analysis_manager )

Collect grid-parallel (GPU-offloaded) ancestor Reduce nodes reducing into container.

Their cross-block merge becomes an atomic writeback, so a localizing transformation demotes them to plain Map nodes and emits an atomic copy-out.

◆ collect_reduction_owners()

bool sdfg::tiles::collect_reduction_owners	(	structured_control_flow::StructuredLoop &	loop,
		const std::string &	container,
		analysis::AnalysisManager &	analysis_manager,
		std::vector< structured_control_flow::Reduce * > &	out )

Collect the non-cooperative Reduce nodes (loop or a descendant) that reduce into container, for accumulator privatization.

Returns: false if a cooperative (GPU-combined) Reduce at/below loop, or a GPU block/warp-cooperative ancestor Reduce, owns container (those are combined by the reduce dispatcher). A sequential or grid-parallel ancestor Reduce is permitted and not added to out. Otherwise fills out with the owning non-cooperative Reduce nodes and returns true.

◆ delinearize_rowmajor()

symbolic::MultiExpression sdfg::tiles::delinearize_rowmajor	(	const symbolic::Expression &	flat,
		const symbolic::MultiExpression &	sizes )

Row-major delinearization of a flat index into per-dim coordinates (dim 0 slowest), matching the buffer's tile linearization so both plan sides agree.

◆ emit_cooperative_copy_loop()

void sdfg::tiles::emit_cooperative_copy_loop	(	codegen::LanguageExtension &	language_extension,
		codegen::PrettyPrinter &	stream,
		const TileCopyNode &	node,
		const std::string &	dst_expr,
		const std::string &	src_expr,
		types::PrimitiveType	elem,
		const std::function< std::string(const std::string &dst_addr, const std::string &src_addr, size_t bytes)> &	async_stmt = {},
		const GuardedLoadFn &	guarded_load = {} )

Emit the block-cooperative staging loop for node into stream: every thread of the enclosing kernel block strides the flat tile from its flat index by the block thread count.

ScalarSync moves one element per step; VectorSync moves a contiguous TileCopyNode::bytes -wide vector per step (the loop strides by the vector's element count). For CpAsync, async_stmt renders the per-step async transfer given the destination/source element addresses and the byte width (the target supplies its own primitive; e.g. CUDA __pipeline_memcpy_async). Guarded (ragged) copy-ins load through guarded_load, defaulting to emit_predicated_guarded_load. Shared by the CUDA and ROCm dispatchers; the visibility barrier is the producer's.

◆ emit_predicated_guarded_load()

std::string sdfg::tiles::emit_predicated_guarded_load	(	const std::string &	type,
		const std::string &	var,
		const std::string &	base,
		const std::string &	index )

Vendor-agnostic GuardedLoadFn: a predicated load into a zero-initialized value (CUDA: @p LDG; AMDGPU: an exec-masked load).

◆ emit_tile_copy_stmt()

std::string sdfg::tiles::emit_tile_copy_stmt	(	codegen::LanguageExtension &	language_extension,
		const TileCopyNode &	node,
		const std::string &	dst_expr,
		const std::string &	src_expr,
		types::PrimitiveType	elem,
		const symbolic::Expression &	index )

The single-element copy statement dst[dst_geo(index)] = src[src_geo(index)]; for the tile coordinate index, addressing both bare-pointer sides as flat elem pointers via the plan's layouts.

Shared by every TileCopyNode dispatcher so the loop shape (sequential / cooperative / async) is the only difference between targets. dst_expr / src_expr are the {_dst,_src} input expressions.

◆ find_block_scheduled_descendant()

structured_control_flow::StructuredLoop * sdfg::tiles::find_block_scheduled_descendant	(	structured_control_flow::StructuredLoop &	loop,
		analysis::AnalysisManager &	analysis_manager )

First block-level GPU-offloaded loop in loop's body (a cooperative copy / consumer axis for enclosing-scope staging), or nullptr.

◆ guard_axis_aligned()

bool sdfg::tiles::guard_axis_aligned	(	const TileCopyNode &	node,
		size_t	axis,
		long long	factor )

Every guard dim on tile axis has a base and exclusive bound that are multiples of factor, so a factor-aligned run of coordinates along it is entirely in or entirely out of bounds.

◆ is_constant_bounded()

bool sdfg::tiles::is_constant_bounded ( const analysis::MemoryTileGroup * group )

True iff group is a real tile whose extents are all compile-time integer constants (the constant-bounded localization precondition).

A null/empty group or any symbolic/unbounded extent is false.

◆ is_reduction_accumulator()

bool sdfg::tiles::is_reduction_accumulator	(	structured_control_flow::StructuredLoop &	loop,
		const std::string &	container,
		analysis::AnalysisManager &	analysis_manager )

Whether container is the accumulator of a Reduce enclosing, nested within, or equal to loop.

◆ localizable_tile()

const analysis::MemoryTileGroup * sdfg::tiles::localizable_tile	(	structured_control_flow::StructuredLoop &	loop,
		const std::string &	container,
		analysis::AnalysisManager &	analysis_manager )

The single localizable tile group of container at loop, or nullptr.

Returns the sole group iff the container forms EXACTLY ONE group there AND every one of its body memlets belongs to it — anchoring on the container so all its access nodes localize together as one coherent tile, which makes wholesale rewriting safe.

◆ provably_multiple_of()

bool sdfg::tiles::provably_multiple_of	(	symbolic::Expression	e,
		long long	m,
		const TileCopyNode &	node )

True if e is a multiple of m for all values of its symbols: the indvar of each loop enclosing node is rewritten as init + stride * t; any other symbol is an unknown integer.

◆ register_tiles_plugin() [1/2]

void sdfg::tiles::register_tiles_plugin ( )

inline

Deprecated: use the variant with explicit context

◆ register_tiles_plugin() [2/2]

void sdfg::tiles::register_tiles_plugin ( plugins::Context & context )

inline

◆ row_pad()

size_t sdfg::tiles::row_pad	(	const BankModel &	model,
		size_t	elem_bytes,
		PackedBuffer	buffer,
		size_t	align_bytes,
		const std::function< std::optional< size_t >(const PackedBuffer &)> &	cost )

The ConflictPolicy::PadRow padding (elements) of buffer minimizing cost, the access cycles of a candidate (nullopt rejects it).

Ties prefer stored rows that stay align_bytes aligned, then less padding. Candidates span one full cycle of the banks.

◆ simplify_tile_bound()

symbolic::Expression sdfg::tiles::simplify_tile_bound	(	const symbolic::Expression &	max,
		const symbolic::Expression &	base,
		const symbolic::Expression &	size )

Simplify a tile's inclusive bound max for coordinates base + i, 0 <= i < size, preserving the predicate base + i <= max on that range: drops min/max arguments dominated by another and min arguments at or above the tile's last coordinate (they never bind).

E.g. max(min(N-1, b+47), min(N-1, b+63)) over a 64-tile at b becomes N-1.

◆ slot_pad()

size_t sdfg::tiles::slot_pad	(	const BankModel &	model,
		size_t	block,
		size_t	coop_span )

ConflictPolicy::PadSlot padding (elements) of a per-slot block of block elements whose slots are accessed by coop_span consecutive lanes of a subgroup: makes the slot stride congruent to coop_span modulo the bank count, rounded up to a multiple of 4 elements so blocks stay vector-addressable.

Without a span, the next odd stride.

◆ tile_element_count()

symbolic::Expression sdfg::tiles::tile_element_count ( const analysis::MemoryTileGroup * group )

Scalar slots a packed buffer for group would occupy (product of extents): a compile-time integer iff is_constant_bounded, null if any extent is unbounded, 0 for a null group.

◆ tile_groups()

const std::vector< analysis::MemoryTileGroup > * sdfg::tiles::tile_groups	(	structured_control_flow::StructuredLoop &	loop,
		const std::string &	container,
		analysis::AnalysisManager &	analysis_manager )

All tile groups MemoryLayoutAnalysis formed for container at loop (nullptr if none).

◆ widest_legal_transfer()

size_t sdfg::tiles::widest_legal_transfer	(	const TileCopyNode &	node,
		CopyAtom	atom,
		const std::vector< size_t > &	widths )

The first of widths (bytes, widest first) at which each lane of node can move one contiguous multi-element run under atom, or 0.

node.bytes() is taken as the element width.

Variable Documentation

◆ LibraryNodeType_PipelineCommit

data_flow::LibraryNodeCode sdfg::tiles::LibraryNodeType_PipelineCommit {"pipeline_commit"}

inline

◆ LibraryNodeType_PipelineWait

data_flow::LibraryNodeCode sdfg::tiles::LibraryNodeType_PipelineWait {"pipeline_wait"}

inline

◆ LibraryNodeType_TileCopy

data_flow::LibraryNodeCode sdfg::tiles::LibraryNodeType_TileCopy {"tile_copy"}

inline

Classes

Typedefs

Enumerations

Functions

Variables

Typedef Documentation

◆ GuardedLoadFn

◆ Layout

Enumeration Type Documentation

◆ ConflictPolicy

◆ CopyAtom

◆ CopyDirection

◆ CopyPhase

◆ Level

◆ Placement

◆ ReductionBufferStatus

◆ Role

◆ Space

Function Documentation

◆ access_cycles()

◆ coalesce_sweep_order()

◆ collect_grid_reduction_owners()

◆ collect_reduction_owners()

◆ delinearize_rowmajor()

◆ emit_cooperative_copy_loop()

◆ emit_predicated_guarded_load()

◆ emit_tile_copy_stmt()

◆ find_block_scheduled_descendant()

◆ guard_axis_aligned()

◆ is_constant_bounded()

◆ is_reduction_accumulator()

◆ localizable_tile()

◆ provably_multiple_of()

◆ register_tiles_plugin() [1/2]

◆ register_tiles_plugin() [2/2]

◆ row_pad()

◆ simplify_tile_bound()

◆ slot_pad()

◆ tile_element_count()

◆ tile_groups()

◆ widest_legal_transfer()

Variable Documentation

◆ LibraryNodeType_PipelineCommit

◆ LibraryNodeType_PipelineWait

◆ LibraryNodeType_TileCopy