Primitive functions that contains TIR statements. More...

#include <function.h>

Inheritance diagram for tvm::tir::PrimFuncNode:

Collaboration diagram for tvm::tir::PrimFuncNode:

Public Member Functions
void	VisitAttrs (tvm::AttrVisitor *v)

bool	SEqualReduce (const PrimFuncNode *other, SEqualReducer equal) const

void	SHashReduce (SHashReducer hash_reduce) const

FuncType	func_type_annotation () const
	Return the derived function annotation of this function. More...

	TVM_OBJECT_ENABLE_SCRIPT_PRINTER ()

	TVM_DECLARE_FINAL_OBJECT_INFO (PrimFuncNode, BaseFuncNode)

Public Member Functions inherited from tvm::BaseFuncNode
template<typename TObjectRef >
Optional< TObjectRef >	GetAttr (const std::string &attr_key, Optional< TObjectRef > default_value=Optional< TObjectRef >(nullptr)) const
	Get a function attribute. More...

template<typename TObjectRef >
Optional< TObjectRef >	GetAttr (const std::string &attr_key, TObjectRef default_value) const

bool	HasNonzeroAttr (const std::string &attr_key) const
	Check whether the function has an non-zero integer attr. More...

LinkageType	GetLinkageType () const
	Get the type of the linkage. More...

	TVM_DECLARE_BASE_OBJECT_INFO (BaseFuncNode, RelayExprNode)

Public Member Functions inherited from tvm::RelayExprNode
const Type &	checked_type () const

template<typename TTypeNode >
const TTypeNode *	type_as () const
	Check if the inferred(checked) type of the Expr is backed by a TTypeNode and return it. More...

VirtualDevice	virtual_device () const

	TVM_DECLARE_BASE_OBJECT_INFO (RelayExprNode, BaseExprNode)

Public Member Functions inherited from tvm::BaseExprNode
	TVM_DECLARE_BASE_OBJECT_INFO (BaseExprNode, 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
Array< tir::Var >	params
	Function parameters. More...

tir::Stmt	body
	The body of the function. More...

Type	ret_type
	The return type of the function. More...

Map< tir::Var, Buffer >	buffer_map
	Maps some parameters to specific Buffer data structures. More...

Public Attributes inherited from tvm::BaseFuncNode
DictAttrs	attrs
	Additional attributes storing the meta-data. More...

Public Attributes inherited from tvm::RelayExprNode
Type	checked_type_ = Type(nullptr)
	Stores the result of type inference(type checking). More...

Optional< ObjectRef >	struct_info_ = Optional<ObjectRef>()
	Stores the result of structure information of the expression that encapsulate both static shape and runtime information such as shape. More...

ObjectRef	virtual_device_
	The virtual device (VirtualDevice) for this node (the result of device planning). For first-order expressions (non functions), this describes where the result of evaluating the expression should be stored. Note that currently, all composite first-order values (tuples, references, ADTs) must be stored on the same virtual device. This means that it is not possible to store two tuple fields on different devices, so we only need one virtual device for these types. More...

Public Attributes inherited from tvm::BaseExprNode
Span	span
	Span that points to the original source code. Reserved debug information. More...

Static Public Attributes
static constexpr const char *	_type_key = "tir.PrimFunc"

Static Public Attributes inherited from tvm::BaseFuncNode
static constexpr const char *	_type_key = "BaseFunc"

static constexpr const uint32_t	_type_child_slots = 2

Static Public Attributes inherited from tvm::RelayExprNode
static constexpr const char *	_type_key = "RelayExpr"

static constexpr const uint32_t	_type_child_slots = 22

Static Public Attributes inherited from tvm::BaseExprNode
static constexpr const char *	_type_key = "BaseExpr"

static constexpr const bool	_type_has_method_sequal_reduce = true

static constexpr const bool	_type_has_method_shash_reduce = true

static constexpr const uint32_t	_type_child_slots = 62

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

Primitive functions that contains TIR statements.

The PrimFunc provides low-level code representation does not automatically manage

See also: PrimFunc

Member Function Documentation

◆ func_type_annotation()

FuncType tvm::tir::PrimFuncNode::func_type_annotation ( ) const

Return the derived function annotation of this function.

Returns: The function type annotation.

Note: The function type annotation of PrimExpr is directly derived from the Vars without the need of type inference.

◆ SEqualReduce()

bool tvm::tir::PrimFuncNode::SEqualReduce	(	const PrimFuncNode *	other,
		SEqualReducer	equal
	)		const

inline

◆ SHashReduce()

void tvm::tir::PrimFuncNode::SHashReduce ( SHashReducer hash_reduce ) const

inline

◆ TVM_DECLARE_FINAL_OBJECT_INFO()

tvm::tir::PrimFuncNode::TVM_DECLARE_FINAL_OBJECT_INFO	(	PrimFuncNode	,
		BaseFuncNode
	)

◆ TVM_OBJECT_ENABLE_SCRIPT_PRINTER()

tvm::tir::PrimFuncNode::TVM_OBJECT_ENABLE_SCRIPT_PRINTER ( )

◆ VisitAttrs()

void tvm::tir::PrimFuncNode::VisitAttrs ( tvm::AttrVisitor * v )

inline

Member Data Documentation

◆ _type_key

constexpr const char* tvm::tir::PrimFuncNode::_type_key = "tir.PrimFunc"

staticconstexpr

◆ body

tir::Stmt tvm::tir::PrimFuncNode::body

The body of the function.

◆ buffer_map

Map<tir::Var, Buffer> tvm::tir::PrimFuncNode::buffer_map

Maps some parameters to specific Buffer data structures.

buffer_map provides a way to express data structure's field and shape constraints. The provided information is used in the program analysis and the code generation.

It defines the vars in the Buffer (m, n) in the cases below when they appears in the buffer_map for the first time.
When a var appears multiple times, they translate into runtime assertion to check the field constraint.

# The corresponding fields of f are as follows
#
# - f.params = [a, b]
# - f.buffer_map = {a: A, b: B}
# - A = decl_buffer(shape=[m, n])
# - B = decl_buffer(shape=[m, n])
 
def f(a, b):
    m, n = var(), var()
    A = bind_buffer(a, shape=[m, n])
    B = bind_buffer(b, shape=[m, n])
    # body

buffer_map is a sugar to express:

Parameter unpacking: e.g. I can load a.shape[0] to get value of m
Constraint checking: a.shape[0] must equal b.shape[0] because they both corresponds to m.

While we could have express parameter unpacking and constraint using normal statements, making buffer_map as first class citizen of PrimFunc will make program analysis much easier.

Prior to buffer flattening, which is performed either in StorageFlatten for TE-based schedules or in FlattenBuffer for TIR-based schedules, these buffer objects are used directly in the body of the function. After buffer flattening, these buffer objects remain unflattened for use in argument validation, but all usage in the body of the function is done through a flattened alias of the buffer.

◆ params

Array<tir::Var> tvm::tir::PrimFuncNode::params

Function parameters.

◆ ret_type

Type tvm::tir::PrimFuncNode::ret_type

The return type of the function.

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

include/tvm/tir/function.h

Public Member Functions

Public Attributes

Static Public Attributes

Additional Inherited Members

Detailed Description

Member Function Documentation

◆ func_type_annotation()

◆ SEqualReduce()

◆ SHashReduce()

◆ TVM_DECLARE_FINAL_OBJECT_INFO()

◆ TVM_OBJECT_ENABLE_SCRIPT_PRINTER()

◆ VisitAttrs()

Member Data Documentation

◆ _type_key

◆ body

◆ buffer_map

◆ params

◆ ret_type