represents a stage. More...

#include <schedule.h>

Inheritance diagram for tvm::te::StageNode:

Collaboration diagram for tvm::te::StageNode:

Public Member Functions
void	VisitAttrs (AttrVisitor *v)

	TVM_DECLARE_FINAL_OBJECT_INFO (StageNode, Object)

Public Member Functions inherited from tvm::runtime::Object
uint32_t	type_index () const

std::string	GetTypeKey () const

size_t	GetTypeKeyHash () const

template<typename TargetType >
bool	IsInstance () const

bool	unique () const

	Object ()

	Object (const Object &other)

	Object (Object &&other)

Object &	operator= (const Object &other)

Object &	operator= (Object &&other)

Public Attributes
Operation	op
	The operation of stage, can be different from original op. If it is null, then this stage is a group stage. More...

Operation	origin_op
	The original operator. The op field can change during schedule to alternate the dataflow, while origin_op remains fixed. More...

Array< IterVar >	all_iter_vars
	All the nodes in the iter var. More...

Array< IterVar >	leaf_iter_vars
	The current active leaf iter vars in the stage. More...

Array< IterVar >	env_threads
	Specify threads to be launched at the stage. This is only valid for composite ops such as Scan. More...

PrimExpr	store_predicate
	The predicate under which store can happen Use this when there can be duplicated threads doing the same store. More...

Array< IterVarRelation >	relations
	The relation bwteen of IterVars. More...

Map< IterVar, IterVarAttr >	iter_var_attrs
	additional attributes about iter var. More...

AttachType	attach_type {kGroupRoot}
	The attachment type of the schedule. More...

IterVar	attach_ivar
	The attach point of this schedule. More...

Stage	attach_stage
	The stage this node attaches to. More...

const ScheduleNode *	attach_sch
	The schedule current stage is attached to. More...

std::string	scope
	The thread storage scope level of the stage. More...

bool	is_output {false}
	Whether this is an output stage. More...

bool	double_buffer {false}
	Whether apply double buffer optimization to this stage. More...

bool	rolling_buffer {false}
	Whether apply rolling buffer optimization to this stage. More...

Array< IndexMap >	layout_transforms
	Layout transformations to be applied onto the stage's tensors. More...

Array< IntImm >	axis_separators
	List of axes after which to divide physical axes. More...

Stage	group
	The parent group of the current stage. The stage cannot be assigned to stages outside the group. More...

int	num_child_stages {0}
	Number of direct child stages, only used for group stage. More...

Static Public Attributes
static constexpr const char *	_type_key = "Stage"

Static Public Attributes inherited from tvm::runtime::Object
static constexpr const char *	_type_key = "runtime.Object"

static constexpr bool	_type_final = false

static constexpr uint32_t	_type_child_slots = 0

static constexpr bool	_type_child_slots_can_overflow = true

static constexpr bool	_type_has_method_visit_attrs = true

static constexpr bool	_type_has_method_sequal_reduce = false

static constexpr bool	_type_has_method_shash_reduce = false

static constexpr uint32_t	_type_index = TypeIndex::kDynamic

Additional Inherited Members
Public Types inherited from tvm::runtime::Object
typedef void(*	FDeleter) (Object *self)
	Object deleter. More...

using	RefCounterType = std::atomic< int32_t >

Static Public Member Functions inherited from tvm::runtime::Object
static std::string	TypeIndex2Key (uint32_t tindex)
	Get the type key of the corresponding index from runtime. More...

static size_t	TypeIndex2KeyHash (uint32_t tindex)
	Get the type key hash of the corresponding index from runtime. More...

static uint32_t	TypeKey2Index (const std::string &key)
	Get the type index of the corresponding key from runtime. More...

static uint32_t	_GetOrAllocRuntimeTypeIndex ()

static uint32_t	RuntimeTypeIndex ()

Protected Member Functions inherited from tvm::runtime::Object
void	IncRef ()
	developer function, increases reference counter. More...

void	DecRef ()
	developer function, decrease reference counter. More...

Static Protected Member Functions inherited from tvm::runtime::Object
static uint32_t	GetOrAllocRuntimeTypeIndex (const std::string &key, uint32_t static_tindex, uint32_t parent_tindex, uint32_t type_child_slots, bool type_child_slots_can_overflow)
	Get the type index using type key. More...

Protected Attributes inherited from tvm::runtime::Object
uint32_t	type_index_ {0}
	Type index(tag) that indicates the type of the object. More...

RefCounterType	ref_counter_ {0}
	The internal reference counter. More...

FDeleter	deleter_ = nullptr
	deleter of this object to enable customized allocation. If the deleter is nullptr, no deletion will be performed. The creator of the object must always set the deleter field properly. More...

Detailed Description

represents a stage.

relations form a Directed acylic hypergraph in bipartite manner. With each node is represented by a IterVar, and each hyper-edge is represented by a IterVarRelation. The relations connects the IterVars in the graph.

Besides typical stage that corresponds to operations. There is also group stage, which groups stages together. Each stage's group(given by group) represent an constraint, the stage can only be attached to stages within the group.

The group stage node can be attached to IterVars as in normal stage.

Member Function Documentation

◆ TVM_DECLARE_FINAL_OBJECT_INFO()

tvm::te::StageNode::TVM_DECLARE_FINAL_OBJECT_INFO	(	StageNode	,
		Object
	)

◆ VisitAttrs()

void tvm::te::StageNode::VisitAttrs ( AttrVisitor * v )

inline

Member Data Documentation

◆ _type_key

constexpr const char* tvm::te::StageNode::_type_key = "Stage"

staticconstexpr

◆ all_iter_vars

Array<IterVar> tvm::te::StageNode::all_iter_vars

All the nodes in the iter var.

Each element of all_iter_vars represents an iteration variable that may appear within this stage's computation. Any element of all_iter_vars that is in leaf_iter_vars represents a variable that is directly defined and usable within the stage's computation. All other elements of all_iter_vars represent variables whose value must be computed from the variables in leaf_iter_vars. (e.g. Support index k has been split by ko, ki = s.split(k, factor=4). ko and ki will appear in leaf_iter_vars, while k will not, and must be computed as 4*ko + ki.

◆ attach_ivar

IterVar tvm::te::StageNode::attach_ivar

The attach point of this schedule.

◆ attach_sch

const ScheduleNode* tvm::te::StageNode::attach_sch

The schedule current stage is attached to.

◆ attach_stage

Stage tvm::te::StageNode::attach_stage

The stage this node attaches to.

◆ attach_type

AttachType tvm::te::StageNode::attach_type {kGroupRoot}

The attachment type of the schedule.

◆ axis_separators

Array<IntImm> tvm::te::StageNode::axis_separators

List of axes after which to divide physical axes.

Used to populate BufferNode::axis_separators, which has additional details.

◆ double_buffer

bool tvm::te::StageNode::double_buffer {false}

Whether apply double buffer optimization to this stage.

◆ env_threads

Array<IterVar> tvm::te::StageNode::env_threads

Specify threads to be launched at the stage. This is only valid for composite ops such as Scan.

Note: Experimental primitive: used for thread persistence.

◆ group

Stage tvm::te::StageNode::group

The parent group of the current stage. The stage cannot be assigned to stages outside the group.

◆ is_output

bool tvm::te::StageNode::is_output {false}

Whether this is an output stage.

◆ iter_var_attrs

Map<IterVar, IterVarAttr> tvm::te::StageNode::iter_var_attrs

additional attributes about iter var.

◆ layout_transforms

Array<IndexMap> tvm::te::StageNode::layout_transforms

Layout transformations to be applied onto the stage's tensors.

◆ leaf_iter_vars

Array<IterVar> tvm::te::StageNode::leaf_iter_vars

The current active leaf iter vars in the stage.

Each element of leaf_iter_vars will either be replaced with the bound index (e.g. threadIdx.x), or will be expanded into a loop over the variable's extent. leaf_iter_vars is a subset of all_iter_vars.

◆ num_child_stages

int tvm::te::StageNode::num_child_stages {0}

Number of direct child stages, only used for group stage.

◆ op

Operation tvm::te::StageNode::op

The operation of stage, can be different from original op. If it is null, then this stage is a group stage.

◆ origin_op

Operation tvm::te::StageNode::origin_op

The original operator. The op field can change during schedule to alternate the dataflow, while origin_op remains fixed.

◆ relations

Array<IterVarRelation> tvm::te::StageNode::relations

The relation bwteen of IterVars.

◆ rolling_buffer

bool tvm::te::StageNode::rolling_buffer {false}

Whether apply rolling buffer optimization to this stage.

◆ scope

std::string tvm::te::StageNode::scope

The thread storage scope level of the stage.

◆ store_predicate

PrimExpr tvm::te::StageNode::store_predicate

The predicate under which store can happen Use this when there can be duplicated threads doing the same store.

Note: Experimental primitive: used by cross thread-reduction.

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

include/tvm/te/schedule.h

Public Member Functions

Public Attributes

Static Public Attributes

Additional Inherited Members