Namespaces
	attr
	PrimFunc specific attribute names.

	builtin
	Collection of builtin intrinsics as ops.

	transform

	usmp

Classes
class	Add
	Managed reference to AddNode. More...

class	AddNode
	a + b More...

class	Allocate
	Managed reference to AllocateNode. More...

class	AllocateConst
	Managed reference to AllocateConstNode. More...

class	AllocateConstNode
	Allocate a buffer that can be used in body. More...

class	AllocateNode
	Allocate a buffer that can be used in body. More...

class	And
	Managed reference to AndNode. More...

class	AndNode
	a && b More...

class	Any
	Managed reference to AnyNode. More...

class	AnyNode
	Any shape. More...

class	AssertStmt
	Managed reference to AssertStmtNode. More...

class	AssertStmtNode
	Assert condition, if an error occurs, return the error message. More...

class	AttrStmt
	Managed reference to AttrStmtNode. More...

class	AttrStmtNode
	Define certain auxiliary attribute for the body to be a symbolic value. This provide auxiliary information for IR passes that transforms body. More...

class	BijectiveLayout
	Bijective function mapping for data layout transformation. Given two Layout, BijectiveLayout build and store the mapping rules, provides API to transform N-dimention tensor from the source indices (i0, i1, .., im) to the destination indices (j0, j1, .., jm). More...

class	BijectiveLayoutNode

class	BinaryOpNode
	Base template to implement binary ops. More...

class	Block
	Managed reference to BlockNode. More...

struct	BlockInfo
	The information about a TensorIR block, it contains two categories of information 1) Info on the block scope rooted at a specific block, including dependency tracking, flags indicating if the scope is a stage pipeline, etc. 2) Info on the block itself, including if the block has a quasi-affine binding, if the regions it reads are completely covered by their producers, etc. More...

class	BlockNode
	A block is a basic schedule unit in TIR. More...

class	BlockRealize
	Managed reference to BlockRealizeNode. More...

class	BlockRealizeNode
	A block realization node represents execution of the block at the binding values. More...

class	BlockRV
	Managed reference to BlockRVNode. More...

class	BlockRVNode
	A random variable that evaluates to a TensorIR block. More...

class	BlockScope
	Managed reference to BlockScopeNode. More...

class	BlockScopeNode
	An object with 1-to-1 correspondence with each block reference in the sref tree. This data structure is used to track the producer-consumer dependencies between blocks. For example even leaf nodes have a scope node, even though they have no dependencies. More...

class	Broadcast
	Managed reference to BroadcastNode. More...

class	BroadcastNode
	Create a vector where all the elements are value. More...

class	Buffer
	Buffer is a symbolic n-darray structure. It is a composition of primitive symbolic types, used to specify the memory layout of the Tensor used in program input. More...

class	BufferLoad
	Managed reference to BufferLoadNode. More...

class	BufferLoadNode
	Load value from the high dimension buffer. More...

class	BufferNode
	Node to represent a buffer. More...

class	BufferRealize
	Managed reference to BufferRealizeNode. More...

class	BufferRealizeNode
	Annotate the region where the buffer need to be read and write in the body. We only need to allocate the space for the corresponding region. More...

class	BufferRegion
	Managed reference to BufferRegionNode. More...

class	BufferRegionNode
	Representing the region of multi-dimensional buffer access. More...

class	BufferStore
	Managed reference to BufferStoreNode. More...

class	BufferStoreNode
	Store value to the high dimension buffer. More...

class	Call
	Managed reference to CallNode. More...

class	CallNode
	Call node. More...

class	Cast
	Managed reference to CastNode. More...

class	CastNode
	Cast value from one data type to another. More...

class	CmpOpNode
	Base template to implement comparison ops. More...

class	CommReducer
	Managed reference to CommReducerNode. More...

class	CommReducerNode
	A commutative reducer node to represent a commutative binary operator with identity element. More...

class	DataProducer
	Managed reference to DataProducerNode. More...

class	DataProducerNode
	Base node for data producers. More...

class	Dependency
	Managed reference to DependencyNode. More...

class	DependencyNode
	A tuple (src, dst, kind) representing certain types of dependency. For example, (A, B, kRAW) means block B depends on block A, and the dependency kind is read-after-write, which means block B reads the result written by block A. More...

class	Div
	Managed reference to DivNode. More...

class	DivNode
	a / b in the C semnatics. More...

class	EQ
	Managed reference to EQNode. More...

class	EQNode
	a == b More...

class	Evaluate
	Managed reference to EvaluateNode. More...

class	EvaluateNode
	Evaluates an expression. This is mostly used for putting a Call node into Stmt. More...

struct	ExprDeepEqual
	Compare two expressions recursively and check if they are equal to each other without var remapping. More...

class	ExprFunctor
	A dynamical functor that dispatches on in the first Expr argument. You can use this as a more powerful Visitor, since it allows you to define function signatures of Visit Function. More...

class	ExprFunctor< R(const PrimExpr &n, Args...)>

class	ExprMutator
	ExprMutator that mutates expressions. More...

class	ExprVisitor
	ExprVisitor. More...

class	FloorDiv
	Managed reference to FloorDivNode. More...

class	FloorDivNode
	Floor division, floor(a/b) More...

class	FloorMod
	Managed reference to FloorModNode. More...

class	FloorModNode
	The remainder of the floordiv. More...

class	For
	Managed reference to ForNode. More...

class	ForNode
	A for loop, with poissible type annotations. More...

class	GE
	Managed reference to GENode. More...

class	GENode
	a >= b More...

class	GT
	Managed reference to GTNode. More...

class	GTNode
	a > b More...

class	IfThenElse
	Managed reference to IfThenElseNode. More...

class	IfThenElseNode
	IfThenElse statment. More...

class	IndexMap

class	IndexMapNode
	Defines a mapping between two representations of indices into a buffer. More...

class	Instruction
	Managed reference to InstructionNode. More...

class	InstructionKind
	Managed reference to InstructionKindNode. More...

class	InstructionKindNode
	Kind of an instruction, e.g. Split, Reorder, etc. Besides the name, every kind of instruction has its own properties, including: 1) A boolean indicating if the instruction is pure, i.e. change nothing in the schedule state 2) A functor that applies the instruction to a TensorIR schedule 3) A functor that converts the instruction to a statement in python syntax 4) A functor that serialize its attributes to JSON 5) A functor that deserialize its attributes from JSON. More...

class	InstructionKindRegEntry
	An entry in the registry of InstructionKind. More...

class	InstructionNode
	Schedule instructions each corresponds to a schedule primitive. More...

class	IterVar
	Iteration Variable, represents an iteration over an integer interval. More...

class	IterVarNode
	An iteration variable representing an iteration over a one dimensional interval. More...

class	Layout
	Managed reference to LayoutNode. More...

class	LayoutAxis

class	LayoutNode
	Layout is to describe how data is organized within an N-dimention tensor. It is composed of upper cases, lower cases and numbers, where upper case indicates a primal axis and the corresponding lower case with factor size indicates the subordinate axis. For example, NCHW16c can describe a 5-D tensor of [batch_size, channel, height, width, channel_block]. Here subordinate axis channel_block=16 is the factor size of the primal axis C (channel). Layout for scalar is defined, while both its name and axes have size 0. More...

class	LE
	Managed reference to LENode. More...

struct	LENode
	a <= b More...

class	Let
	Managed reference to LetNode. More...

class	LetNode
	Let binding. Bind var to value then evaluate body. More...

class	LetStmt
	Managed reference to LetStmtNode. More...

class	LetStmtNode
	Let binding, bind var to value, then run body. More...

class	Load
	Managed reference to LoadNode. More...

class	LoadNode
	Load the value from buffer_var. More...

class	LoopRV
	Managed reference to LoopRVNode. More...

class	LoopRVNode
	A random variable that evaluates to a TensorIR for loop. More...

class	LT
	Managed reference to LTNode. More...

class	LTNode
	a < b More...

class	MatchBufferRegion
	Managed reference to MatchBufferRegionNode. More...

class	MatchBufferRegionNode
	Match introduces a constraint that the source buffer region can be remapped to the data layout specified by the buffer field. The constraint can be checked in later part of lowering (or optionally during runtime). More...

class	Max
	Managed reference to MaxNode. More...

class	MaxNode
	max(a, b) More...

class	Min
	Managed reference to MinNode. More...

class	MinNode
	min(a, b) More...

class	Mod
	Managed reference to ModNode. More...

class	ModNode
	a % b in the C semnatics. More...

class	Mul
	Managed reference to MulNode. More...

class	MulNode
	a * b More...

class	NE
	Managed reference to NENode. More...

class	NENode
	a != b More...

class	Not
	Managed reference to NotNode. More...

class	NotNode
	!a More...

class	Or
	Managed reference to OrNode. More...

class	OrNode
	a \|\| b More...

class	Prefetch
	Managed reference to PrefetchNode. More...

class	PrefetchNode
	A prefetch hint for a buffer. More...

class	PrimFunc
	Managed reference to PrimFuncNode. More...

class	PrimFuncNode
	Primitive functions that contains TIR statements. More...

class	ProducerLoad
	Managed reference to ProducerLoadNode. More...

class	ProducerLoadNode
	Load value from the result produced by the producer. More...

class	ProducerRealize
	Managed reference to ProducerRealizeNode. More...

class	ProducerRealizeNode
	Annotate the bounds where the data produced by the producer need to be written and read in body. We will need to allocate space for the corresponding regions. More...

class	ProducerStore
	Managed reference to ProducerStoreNode. More...

class	ProducerStoreNode
	Store value into mult-dimensional array that will be read by the consumer of the producer. More...

class	Ramp
	Managed reference to RampNode. More...

class	RampNode
	Construct a vector with lanes elements where its i-th element equals base + i * stride. This is useful to construct a index for a continuous vector load. More...

class	Reduce
	Managed reference to ReduceNode. More...

class	ReduceNode
	Reduction operator operator. More...

class	Schedule
	Managed reference to ScheduleNode. More...

class	ScheduleNode
	The user-facing schedule class. More...

class	ScheduleState
	Managed reference to ScheduleStateNode. More...

class	ScheduleStateNode
	The state of scheduling, which exposes a `Replace` method as the primary interface for all the scheduling primitives to manipulate the TensorIR. More...

class	Select
	Managed reference to SelectNode. More...

class	SelectNode
	return true_value if condition is true, otherwise return false_value. More...

class	SeqStmt
	Sequence statement. More...

class	SeqStmtNode
	The container of seq statement. Represent a sequence of statements. More...

class	Shuffle
	Managed reference to ShuffleNode. More...

class	ShuffleNode
	Shuffle instruction. vec = concat(vectors) result = (vec[indices[0]], vec[indices[1]] ...) More...

class	SizeVar
	a named variable represents a tensor index size More...

class	SizeVarNode
	A variable node represent a tensor index size, whose value must be non-negative. More...

class	Stmt
	Container of all statements. More...

class	StmtExprMutator
	Mutator that recursively mutates stmts and exprs on them. More...

class	StmtExprVisitor
	Visitor that recursively visit stmts and exprs on them. More...

class	StmtFunctor
	Same as ExprFunctor except it is applied on statements. More...

class	StmtFunctor< R(const Stmt &n, Args... args)>

class	StmtMutator
	StmtMutator that mutates the statements. More...

class	StmtNode
	Base node of all statements. More...

class	StmtSRef
	Managed reference to StmtSRefNode. More...

class	StmtSRefNode
	An object that refers to schedulable elements (block/for-loop) in TensorIR, aka "sref". More...

class	StmtVisitor
	StmtVisitor. More...

class	Store
	Managed reference to StoreNode. More...

class	StoreNode
	Store value to the buffer. More...

class	StringImm
	Managed reference to StringImmNode. More...

class	StringImmNode
	String constants, only used in asserts. More...

class	Sub
	Managed reference to SubNode. More...

class	SubNode
	a - b More...

class	TensorIntrin
	Managed reference to TensorIntrinNode. More...

class	TensorIntrinNode
	Tensor intrinsics for tensorization. More...

class	Trace
	Managed reference to TraceNode. More...

class	TraceNode
	An execution trace of a scheduling program. More...

class	Var
	a named variable in TIR More...

class	VarNode
	A variable node in the IR. More...

class	While
	Managed reference to WhileNode. More...

class	WhileNode
	A While loop. More...

Typedefs
using	IntImmNode = tvm::IntImmNode

using	FloatImmNode = tvm::FloatImmNode

using	TGlobalSymbol = String
	Global symbol of the op after lowering. More...

using	TVectorizable = bool
	Whether the op is overloaded for vector form. More...

using	FLowerIntrinsic = runtime::TypedPackedFunc< PrimExpr(PrimExpr)>
	The intrinsic lowering function for given op. More...

using	FLegalize = runtime::TypedPackedFunc< PrimExpr(PrimExpr)>
	The legalization function for given tir op. More...

using	TCallEffectKind = Integer
	Use integer to record the kind. More...

using	FInstructionApply = runtime::TypedPackedFunc< Array< ObjectRef >(Schedule sch, const Array< ObjectRef > &inputs, const Array< ObjectRef > &attrs, const Optional< ObjectRef > &decision)>
	Type of the functor that applies the instruction to a TensorIR schedule. More...

using	FInstructionAsPython = runtime::TypedPackedFunc< String(const Array< ObjectRef > &inputs, const Array< ObjectRef > &attrs, const Optional< ObjectRef > &decision, const Array< String > &outputs)>
	Type of the functor that converts the instruction to a statement in python syntax. More...

using	FInstructionAttrsAsJSON = runtime::TypedPackedFunc< ObjectRef(Array< ObjectRef > attrs)>
	Type of the functor that serialize its attributes to JSON. More...

using	FInstructionAttrsFromJSON = runtime::TypedPackedFunc< Array< ObjectRef >(ObjectRef json_attrs)>
	Type of the functor that deserialize its attributes from JSON. More...

using	ExprRV = PrimExpr
	An expr random variable. More...

using	ExprRVNode = PrimExprNode

using	FTraceDecisionProvider = runtime::TypedPackedFunc< ObjectRef(const Instruction &inst, const Array< ObjectRef > &inputs, const Array< ObjectRef > &attrs, const Optional< ObjectRef > &decision)>
	A callback that allows users to mutate decisions on the fly when applying instructions. The signature of the callback is: More...

using	Region = Array< Range >

Enumerations
enum	BufferType : int { kDefault = 1, kAutoBroadcast = 2 }
	buffer type More...

enum	CallEffectKind : int { CallEffectKind::kExprAnnotation = 0, CallEffectKind::kPure = 1, CallEffectKind::kReadState = 2, CallEffectKind::kUpdateState = 3, kOpaque = kUpdateState, CallEffectKind::kSpecialCallArg = 4, CallEffectKind::kEmbedInfo = 5, CallEffectKind::kControlJump = 6 }
	The effect type of the call. More...

enum	DepKind : int32_t { DepKind::kRAW = 0, DepKind::kWAW = 1, DepKind::kWAR = 2, kOpaque = 3 }
	Type of dependency. Right now we have 4 types of dependencies 1) Read-after-write (kRAW) 2) Write-after-write (kWAW) 3) Write-after-read (kWAR) 4) Opaque dependency (kOpaque) More...

enum	ScheduleErrorRenderLevel : int32_t { ScheduleErrorRenderLevel::kDetail = 0, ScheduleErrorRenderLevel::kFast = 1, ScheduleErrorRenderLevel::kNone = 2 }
	The level of detailed error message rendering. More...

enum	BufferIndexType : int32_t { BufferIndexType::kRead = 0, BufferIndexType::kWrite = 1 }
	Type of buffer index. More...

enum	ScheduleDebugMask : uint32_t { kVerifySRefTree = 1, kVerifyCachedFlags = 2 }
	The bitmask of the debug flag in the ScheduleStateNode. More...

enum	ForKind : int { ForKind::kSerial = 0, ForKind::kParallel = 1, kVectorized = 2, kUnrolled = 3, ForKind::kThreadBinding = 4 }
	The kind of the loop. More...

enum	IterVarType : int { kDataPar = 0, kThreadIndex = 1, kCommReduce = 2, kOrdered = 3, kOpaque = 4, kOpaque = kUpdateState, kOpaque = 3, kUnrolled = 5, kUnrolled = 3, kVectorized = 6, kVectorized = 2, kParallelized = 7, kTensorized = 8 }
	Type of iteration variable. Each IterVar have a specific type. More...

Functions
template<class FLambda >
void	VisitPrimFuncs (const IRModule &mod, FLambda fvisit)
	Visit the PrimFuncs in the IRModule. More...

double	EstimateTIRFlops (const Stmt &stmt)
	Estimate the FLOPs of a TIR fragment. More...

double	EstimateTIRFlops (const IRModule &mod)
	Estimate the FLOPs of TIRs in an IRModule. More...

Array< Var >	UndefinedVars (const Stmt &stmt, const Array< Var > &defs)
	Find undefined vars in the statement. More...

Array< Var >	UndefinedVars (const PrimExpr &expr)
	Find undefined vars in the expression. More...

CallEffectKind	SideEffect (const PrimExpr &expr)
	Analyze the side effect. More...

bool	UsesVar (const Stmt &stmt, std::function< bool(const VarNode *)> vset_contains)
	Whether the given Stmt uses any var in the given variable set. More...

bool	UsesVar (const PrimExpr &expr, std::function< bool(const VarNode *)> vset_contains)
	Whether the given PrimExpr uses any var in the given variable set. More...

bool	VerifySSA (const PrimFunc &func)
	Verifies whether the IR stmt or Expr is in SSA form. That is: each Var is defined and assigned once(in Let/For) More...

bool	VerifyMemory (const PrimFunc &func)
	Verify if memory accesses are legal for a specific target device type. More...

bool	VerifyGPUCode (const PrimFunc &func, Map< String, PrimExpr > constraints)
	Verify the correctness of a GPU code It will check the whether the amount of memory usage or the number of threads in a block exceeds the limit. More...

Array< Array< BufferRegion > >	GetBlockAccessRegion (const Block &block, const Map< Var, Buffer > &buffer_var_map)
	Auto detect the block access region according to its body stmt It will detect the access region as an array in order of appearance in AST. More...

Array< Array< BufferRegion > >	GetBlockReadWriteRegion (const Block &block, const Map< Var, Buffer > &buffer_var_map)
	Auto detect the block read/write region according to its body stmt. An opaque access will be counted as both a read and a write access. More...

size_t	CalculateExprComplexity (const PrimExpr &expr)
	Calculate the expresion complexity based on number of symbols it contains. More...

size_t	CalculateWorkspaceBytes (const PrimFunc &func, const Integer &workspace_byte_alignment)
	Calculate the workspace size in bytes needed by the TIR allocates inside the TIR PrimFunc. More...

Map< Buffer, Optional< Stmt > >	DetectBufferAccessLCA (const PrimFunc &func)
	Detect the lowest common ancestor(LCA) of buffer access, including both high-level access(BufferLoad, BufferStore) and low-level access(Load, Store and opaque access). The LCA may be a For loop or a Block. More...

Buffer	decl_buffer (Array< PrimExpr > shape, DataType dtype=DataType::Float(32), String name="buffer", String storage_scope="", Array< IntImm > axis_separators={}, Span span=Span())
	Construct a new buffer given shape, and dtype. More...

template<typename K , typename V >
std::unordered_map< K, V >	as_unordered_map (const Map< K, V > &dmap)

PrimFunc	Specialize (PrimFunc func, const Map< Var, ObjectRef > &param_map)

bool	IsPointerType (const Type &type, const DataType &element_type)
	Check if type is a pointer to a runtime element type. More...

template<typename ValueType , typename = typename std::enable_if<std::is_pod<ValueType>::value>::type>
PrimExpr	make_const (DataType t, ValueType value, Span span=Span())
	Make a const value with certain data type. More...

PrimExpr	make_zero (DataType t, Span span=Span())
	Make a const zero expr. More...

PrimExpr	const_true (int lanes=1, Span span=Span())
	Make a constant true expression. More...

PrimExpr	const_false (int lanes=1, Span span=Span())
	Make a constant false expression. More...

const int64_t *	as_const_int (const PrimExpr &x)
	Get x as constant int expression. More...

bool	is_const_int (const PrimExpr &x, int64_t value)
	Check whether x is a constant integer expression. More...

bool	is_no_op (const tir::Stmt &stmt)
	Check whether stmt is nop. More...

bool	is_one (const PrimExpr &x)
	Check whether x is a constant integer 1. More...

bool	is_zero (const PrimExpr &x)
	Check whether x is a constant integer 0. More...

bool	is_const_int (const PrimExpr &x)
	Check whether x is an integer constant. More...

bool	is_const_number (const PrimExpr &x)
	Check whether x is an integer/float constant. More...

template<typename FReduce >
PrimExpr	foldl (FReduce freduce, PrimExpr init_value, const Array< PrimExpr > &values, Span span=Span())
	Left fold. More...

bool	is_const_power_of_two_integer (const PrimExpr &x, int *shift)
	Check whether x is a constant power of two If x is power of two, write the power to the shift. More...

bool	is_positive_const (const PrimExpr &a)

bool	is_negative_const (const PrimExpr &a)

template<typename ValueType >
PrimExpr	MakeConstScalar (DataType t, ValueType value, Span span=Span())

PrimExpr	TypeAnnotation (DataType dtype, Span span=Span())
	Create a type annotation expression. More...

std::ostream &	operator<< (std::ostream &os, ForKind kind)

const char *	ForKind2String (ForKind t)

Stmt	IRTransform (Stmt stmt, const runtime::PackedFunc &preorder, const runtime::PackedFunc &postorder, Optional< Array< String >> only_enable=NullOpt)
	recursively visit the ir nodes in post DFS order, and transform it More...

void	PostOrderVisit (const ObjectRef &node, std::function< void(const ObjectRef &)> fvisit)
	Recursively visit the ir in post DFS order node, apply fvisit Each node is guaranteed to be visited only once. More...

Stmt	Substitute (Stmt stmt, std::function< Optional< PrimExpr >(const Var &var)> vmap)
	Substitute the var specified by vmap. More...

PrimExpr	Substitute (PrimExpr expr, std::function< Optional< PrimExpr >(const Var &var)> vmap)
	Substitute the var specified by vmap. More...

Array< Range >	Substitute (const Array< Range > &region, const Map< Var, PrimExpr > &vmap)
	Substitute the var specified by vmap. More...

template<typename T >
auto	Substitute (T input, const Map< Var, PrimExpr > &value_map)
	Sugar for substitute via a given map. More...

template<typename T >
T	Substitute (T input, const std::unordered_map< const VarNode *, PrimExpr > &value_map)
	Sugar for substitute via a given map. More...

void	PreOrderVisit (const ObjectRef &stmt_or_expr, const std::function< bool(const ObjectRef &)> &fvisit)
	Recursively visit the IR in pre DFS order node, apply fvisit. If fvisit returns false, it won't visit the children of the node. More...

PrimFunc	RenewDefs (const PrimFunc &func)
	Renew the definition nodes for a TIR, including Var, Buffer and IterVar. This pass works as a simple DeepCopy to duplicate a function with different Vars and Buffers but the same behavior. More...

const char *	IterVarType2String (IterVarType t)

Typedef Documentation

◆ ExprRV

using tvm::tir::ExprRV = typedef PrimExpr

An expr random variable.

◆ ExprRVNode

using tvm::tir::ExprRVNode = typedef PrimExprNode

◆ FInstructionApply

using tvm::tir::FInstructionApply = typedef runtime::TypedPackedFunc<Array<ObjectRef>( Schedule sch, const Array<ObjectRef>& inputs, const Array<ObjectRef>& attrs, const Optional<ObjectRef>& decision)>

Type of the functor that applies the instruction to a TensorIR schedule.

Parameters

sch	The schedule to be applied on
inputs	The input random variables
attrs	Instruction attributes
decision	Decisions made on the instruction

Returns: The functor returns an array of output random variables

◆ FInstructionAsPython

using tvm::tir::FInstructionAsPython = typedef runtime::TypedPackedFunc<String( const Array<ObjectRef>& inputs, const Array<ObjectRef>& attrs, const Optional<ObjectRef>& decision, const Array<String>& outputs)>

Type of the functor that converts the instruction to a statement in python syntax.

Parameters

inputs	Names of the input random variables
attrs	Instruction attributes
decisions	Decisions made on the instruction
outputs	Names of the output random variables

Returns: A string representing the python api call

◆ FInstructionAttrsAsJSON

using tvm::tir::FInstructionAttrsAsJSON = typedef runtime::TypedPackedFunc<ObjectRef(Array<ObjectRef> attrs)>

Type of the functor that serialize its attributes to JSON.

Parameters

attrs The attributes to be serialized

Returns: An array, serialized attributes

Note: This functor is nullable

◆ FInstructionAttrsFromJSON

using tvm::tir::FInstructionAttrsFromJSON = typedef runtime::TypedPackedFunc<Array<ObjectRef>(ObjectRef json_attrs)>

Type of the functor that deserialize its attributes from JSON.

Parameters

json_attrs The attributes to be serialized

Returns: An array, deserialized attributes

Note: This functor is nullable

◆ FLegalize

using tvm::tir::FLegalize = typedef runtime::TypedPackedFunc<PrimExpr(PrimExpr)>

The legalization function for given tir op.

◆ FloatImmNode

using tvm::tir::FloatImmNode = typedef tvm::FloatImmNode

◆ FLowerIntrinsic

using tvm::tir::FLowerIntrinsic = typedef runtime::TypedPackedFunc<PrimExpr(PrimExpr)>

The intrinsic lowering function for given op.

◆ FTraceDecisionProvider

using tvm::tir::FTraceDecisionProvider = typedef runtime::TypedPackedFunc<ObjectRef( const Instruction& inst, const Array<ObjectRef>& inputs, const Array<ObjectRef>& attrs, const Optional<ObjectRef>& decision)>

A callback that allows users to mutate decisions on the fly when applying instructions. The signature of the callback is:

Parameters

inst	The instruction
inputs	The input random variables
attrs	The attributes
decision	The original decision

Returns: A new decision

◆ IntImmNode

using tvm::tir::IntImmNode = typedef tvm::IntImmNode

◆ Region

using tvm::tir::Region = typedef Array<Range>

◆ TCallEffectKind

using tvm::tir::TCallEffectKind = typedef Integer

Use integer to record the kind.

◆ TGlobalSymbol

using tvm::tir::TGlobalSymbol = typedef String

Global symbol of the op after lowering.

◆ TVectorizable

using tvm::tir::TVectorizable = typedef bool

Whether the op is overloaded for vector form.

Enumeration Type Documentation

◆ BufferIndexType

enum tvm::tir::BufferIndexType : int32_t

strong

Type of buffer index.

Enumerator
kRead	Index of a read buffer.
kWrite	Index of a written buffer.

◆ BufferType

enum tvm::tir::BufferType : int

buffer type

Enumerator
kDefault
kAutoBroadcast

◆ CallEffectKind

enum tvm::tir::CallEffectKind : int

strong

The effect type of the call.

Enumerator
kExprAnnotation	Function corresponds to an annotation(e.g. likely) and can translate to identity.
kPure	Pure function that do not interacts with any external state.
kReadState	Function's that may read from states(e.g. RAM)
kUpdateState	Function that may read/write from states(e.g. RAM).
kOpaque	Opaque function, cannot make any assumption.
kSpecialCallArg	Special intrinsic to annotate call arguments info only valid as a direct argument to a call.
kEmbedInfo	Embed opaque information in the Expr, cannot be codegen.
kControlJump	Function that changes control flow.

◆ DepKind

enum tvm::tir::DepKind : int32_t

strong

Type of dependency. Right now we have 4 types of dependencies 1) Read-after-write (kRAW) 2) Write-after-write (kWAW) 3) Write-after-read (kWAR) 4) Opaque dependency (kOpaque)

Enumerator
kRAW
kWAW
kWAR
kOpaque

◆ ForKind

enum tvm::tir::ForKind : int

strong

The kind of the loop.

ForKind can change the control flow semantics of the loop. So the kind field needs to be considered in all TIR passes.

Enumerator
kSerial	default semantics – serial execution.
kParallel	Parallel execution on CPU.
kVectorized	Vector SIMD loop. The loop body will be vectorized.
kUnrolled	The loop body must be unrolled.
kThreadBinding	The loop variable is bound to a thread in an environment. In the final stage of lowering, the loop is simply removed and the loop variable is mapped to the corresponding context thread.

◆ IterVarType

enum tvm::tir::IterVarType : int

Type of iteration variable. Each IterVar have a specific type.

The type of iter var can be overriden via stage.iter_var_attrs given they are compatible.

Enumerator
kDataPar	Data parallel iteration. This normally corresponds to axis of Tensor. Allow all IterVar manipulations. Note This does not mean the loop have to be executed in parallel fashion.
kThreadIndex	The IterVar itself is a thread-index of a fixed thread launching group. Note that this is already assumed to be parallelized. Disallow: split/fuse/vectorize/parallel
kCommReduce	Communicative reduction. Cannot be directly parallelized. Disallow: parallel/vectorize
kOrdered	Serial loops with loop carry dependency, the iteration must execute in order. Cannot be re-ordered. Disallow: reorder/parallel/vectorize
kOpaque	IterVar is opaque,. May not corresponds to any generated loop Disallow all IterVar manipulations and compute_at Note This is usually used to implement composite op or external op, where the
kOpaque	Opaque function, cannot make any assumption.
kOpaque
kUnrolled	The execution is unrolled.
kUnrolled	The loop body must be unrolled.
kVectorized	The loop is vectorized.
kVectorized	Vector SIMD loop. The loop body will be vectorized.
kParallelized	The loop is parallelized.
kTensorized	Marks boundary of tensorization intrinsic.

◆ ScheduleDebugMask

enum tvm::tir::ScheduleDebugMask : uint32_t

The bitmask of the debug flag in the ScheduleStateNode.

See also: ScheduleStateNode

Enumerator
kVerifySRefTree	Verify the correctness of the sref tree.
kVerifyCachedFlags	Verify the correctness of affine_binding, region_cover and stage_pipeline.

◆ ScheduleErrorRenderLevel

enum tvm::tir::ScheduleErrorRenderLevel : int32_t

strong

The level of detailed error message rendering.

Enumerator
kDetail	Render a detailed error message.
kFast	Render the error in fast mode.
kNone	No error message at all.

Function Documentation

◆ as_const_int()

const int64_t* tvm::tir::as_const_int ( const PrimExpr & x )

inline

Get x as constant int expression.

Parameters

x	The expression

Returns: the address to the int expression, return nullptr, if x is not IntImm.

◆ as_unordered_map()

template<typename K , typename V >

std::unordered_map<K, V> tvm::tir::as_unordered_map ( const Map< K, V > & dmap )

inline

◆ CalculateExprComplexity()

size_t tvm::tir::CalculateExprComplexity ( const PrimExpr & expr )

Calculate the expresion complexity based on number of symbols it contains.

Parameters

expr	The expr to be calculated.

◆ CalculateWorkspaceBytes()

size_t tvm::tir::CalculateWorkspaceBytes	(	const PrimFunc &	func,
		const Integer &	workspace_byte_alignment
	)

Calculate the workspace size in bytes needed by the TIR allocates inside the TIR PrimFunc.

Parameters

func	The TIR PrimFunc for which the workspace size to be calculated
workspace_byte_alignment	The byte alignment required for each tensor allocated in this workspace

◆ const_false()

PrimExpr tvm::tir::const_false	(	int	lanes = `1`,
		Span	span = `Span()`
	)

inline

Make a constant false expression.

Parameters

lanes	The number of lanes in the bool
span	The location of this operation in the source.

Returns: The result expression.

◆ const_true()

PrimExpr tvm::tir::const_true	(	int	lanes = `1`,
		Span	span = `Span()`
	)

inline

Make a constant true expression.

Parameters

lanes	The number of lanes in the bool
span	The location of this operation in the source.

Returns: The result expression.

◆ decl_buffer()

Buffer tvm::tir::decl_buffer	(	Array< PrimExpr >	shape,
		DataType	dtype = `DataType::Float(32)`,
		String	name = `"buffer"`,
		String	storage_scope = `""`,
		Array< IntImm >	axis_separators = `{}`,
		Span	span = `Span()`
	)

Construct a new buffer given shape, and dtype.

Parameters

shape	The shape of the buffer,
dtype	The content data type.
name	The name of the buffer
storage_scope	The storage scope associated with this buffer
axis_separators	Divisions defining the groups of axes that will be flattened together.
span	The location of this object in the source code.

Returns: The created buffer.

See also: Buffer for complete constructor.

◆ DetectBufferAccessLCA()

Map<Buffer, Optional<Stmt> > tvm::tir::DetectBufferAccessLCA ( const PrimFunc & func )

Detect the lowest common ancestor(LCA) of buffer access, including both high-level access(BufferLoad, BufferStore) and low-level access(Load, Store and opaque access). The LCA may be a For loop or a Block.

Parameters

func	The PrimFunc to be detected.

Returns: The Map from buffer to the LCA of all access to it. The lca is function root if the return stmt is NullOpt.

◆ EstimateTIRFlops() [1/2]

double tvm::tir::EstimateTIRFlops ( const Stmt & stmt )

Estimate the FLOPs of a TIR fragment.

Parameters

stmt	The TIR fragment to be estimated.

Returns: The estimated FLOPs.

◆ EstimateTIRFlops() [2/2]

double tvm::tir::EstimateTIRFlops ( const IRModule & mod )

Estimate the FLOPs of TIRs in an IRModule.

Parameters

mod	The IRModule to be estimated.

Returns: The estimated FLOPs.

◆ foldl()

template<typename FReduce >

PrimExpr tvm::tir::foldl	(	FReduce	freduce,
		PrimExpr	init_value,
		const Array< PrimExpr > &	values,
		Span	span = `Span()`
	)

inline

Left fold.

Parameters

freduce	The reduction function.
init_value	The initial value.
values	The values to be folded.
span	The location of the fold in the source.

Returns: The result.

Template Parameters

FReduce The type of the reduction.

◆ ForKind2String()

const char* tvm::tir::ForKind2String ( ForKind t )

inline

◆ GetBlockAccessRegion()

Array<Array<BufferRegion> > tvm::tir::GetBlockAccessRegion	(	const Block &	block,
		const Map< Var, Buffer > &	buffer_var_map
	)

Auto detect the block access region according to its body stmt It will detect the access region as an array in order of appearance in AST.

Parameters

block	The block to be detected
buffer_var_map	The outside buffers which may be accessed the block. It is a map from buffer var to the buffer.

Returns

Array of access regions. There are three arrays of BufferRegion:

first: read regions
second: write regions
third: opaque regions

◆ GetBlockReadWriteRegion()

Array<Array<BufferRegion> > tvm::tir::GetBlockReadWriteRegion	(	const Block &	block,
		const Map< Var, Buffer > &	buffer_var_map
	)

Auto detect the block read/write region according to its body stmt. An opaque access will be counted as both a read and a write access.

Parameters

block	The block to be detected
buffer_var_map	The outside buffers which may be accessed the block. It is a map from buffer var to the buffer

Returns: An array only consisting of the read regions and write regions of the input block

◆ IRTransform()

Stmt tvm::tir::IRTransform	(	Stmt	stmt,
		const runtime::PackedFunc &	preorder,
		const runtime::PackedFunc &	postorder,
		Optional< Array< String >>	only_enable = `NullOpt`
	)

recursively visit the ir nodes in post DFS order, and transform it

Parameters

stmt	The ir to be transformed.
preorder	The function called in before recursive mutation If preorder returns None, then the transform will proceed to recursive call. If preorder returns a not None Stmt/Expr, the transformer will simply return it and won't do further recursion.
postorder	The function called after recursive mutation. The recursive mutation result is passed to postorder for further mutation.
only_enable	List of runtime::String. If it is null, all IRNode will call preorder/postorder If it is not null, preorder/postorder will only be called when the IRNode's type key is in the list.

◆ is_const_int() [1/2]

bool tvm::tir::is_const_int	(	const PrimExpr &	x,
		int64_t	value
	)

inline

Check whether x is a constant integer expression.

Parameters

x	The input argument
value	the value to be compared against.

Returns: whether x is constant expression.

◆ is_const_int() [2/2]

bool tvm::tir::is_const_int ( const PrimExpr & x )

inline

Check whether x is an integer constant.

Note: This only return true for integer types.

Returns: whether x is constant

◆ is_const_number()

bool tvm::tir::is_const_number ( const PrimExpr & x )

inline

Check whether x is an integer/float constant.

Note: This only return true for integer types.

Returns: whether x is constant

◆ is_const_power_of_two_integer()

bool tvm::tir::is_const_power_of_two_integer	(	const PrimExpr &	x,
		int *	shift
	)

Check whether x is a constant power of two If x is power of two, write the power to the shift.

Parameters

x	The input expression.
shift	The output shift if x is power of two.

Returns: whether x is constant power of two

◆ is_negative_const()

bool tvm::tir::is_negative_const ( const PrimExpr & a )

inline

◆ is_no_op()

bool tvm::tir::is_no_op ( const tir::Stmt & stmt )

inline

Check whether stmt is nop.

Parameters

stmt	The input statement

Returns: whether stmt is nop

◆ is_one()

bool tvm::tir::is_one ( const PrimExpr & x )

inline

Check whether x is a constant integer 1.

Parameters

x	The input argument.

Note: This only return true for integer types.

Returns: whether x is constant 1

◆ is_positive_const()

bool tvm::tir::is_positive_const ( const PrimExpr & a )

inline

◆ is_zero()

bool tvm::tir::is_zero ( const PrimExpr & x )

inline

Check whether x is a constant integer 0.

Parameters

x	The input argument

Returns: whether x is constant 0

Note: This only return true for integer types.

◆ IsPointerType()

bool tvm::tir::IsPointerType	(	const Type &	type,
		const DataType &	element_type
	)

inline

Check if type is a pointer to a runtime element type.

Parameters

type	The type to be checked.
element_type	The corresponding element type.

Returns: The check results

◆ IterVarType2String()

const char* tvm::tir::IterVarType2String ( IterVarType t )

inline

◆ make_const()

template<typename ValueType , typename = typename std::enable_if<std::is_pod<ValueType>::value>::type>

PrimExpr tvm::tir::make_const	(	DataType	t,
		ValueType	value,
		Span	span = `Span()`
	)

inline

Make a const value with certain data type.

Parameters

t	The target type.
value	The input value

Returns: the result expression.

Template Parameters

ValueType The constant value type

Parameters

span	The location of this operation in the source.

◆ make_zero()

PrimExpr tvm::tir::make_zero	(	DataType	t,
		Span	span = `Span()`
	)

inline

Make a const zero expr.

Parameters

t	The target type.
span	The location of this operation in the source.

Returns: the result expression.

◆ MakeConstScalar()

template<typename ValueType >

PrimExpr tvm::tir::MakeConstScalar	(	DataType	t,
		ValueType	value,
		Span	span = `Span()`
	)

inline

◆ operator<<()

std::ostream& tvm::tir::operator<<	(	std::ostream &	os,
		ForKind	kind
	)

◆ PostOrderVisit()

void tvm::tir::PostOrderVisit	(	const ObjectRef &	node,
		std::function< void(const ObjectRef &)>	fvisit
	)

Recursively visit the ir in post DFS order node, apply fvisit Each node is guaranteed to be visited only once.

Parameters

node	The ir to be visited.
fvisit	The visitor function to be applied.

◆ PreOrderVisit()

void tvm::tir::PreOrderVisit	(	const ObjectRef &	stmt_or_expr,
		const std::function< bool(const ObjectRef &)> &	fvisit
	)

Recursively visit the IR in pre DFS order node, apply fvisit. If fvisit returns false, it won't visit the children of the node.

Parameters

stmt_or_expr	The ir to be visited.
fvisit	The visitor function to be applied. If fvisit returns false, it won't visit the children of the node

◆ RenewDefs()

PrimFunc tvm::tir::RenewDefs ( const PrimFunc & func )

Renew the definition nodes for a TIR, including Var, Buffer and IterVar. This pass works as a simple DeepCopy to duplicate a function with different Vars and Buffers but the same behavior.

Parameters

func	The input PrimFunc.

Returns: The renewed func.

◆ SideEffect()

CallEffectKind tvm::tir::SideEffect ( const PrimExpr & expr )

Analyze the side effect.

Parameters

expr	The expression to be checked.

Returns: CallEffectKind, can be kPure, kReadState or kUpdateState

◆ Specialize()

PrimFunc tvm::tir::Specialize	(	PrimFunc	func,
		const Map< Var, ObjectRef > &	param_map
	)

◆ Substitute() [1/5]

Stmt tvm::tir::Substitute	(	Stmt	stmt,
		std::function< Optional< PrimExpr >(const Var &var)>	vmap
	)

Substitute the var specified by vmap.

Parameters

stmt	The source statement to be substituted
vmap	returns a new value if re-mapping is needed, otherwise returns nullptr.

Returns: The converted form.

◆ Substitute() [2/5]

PrimExpr tvm::tir::Substitute	(	PrimExpr	expr,
		std::function< Optional< PrimExpr >(const Var &var)>	vmap
	)

Substitute the var specified by vmap.

Parameters

expr	The source statement to be substituted
vmap	returns a new value if re-mapping is needed, otherwise returns nullptr.

Returns: The result.

◆ Substitute() [3/5]

Array<Range> tvm::tir::Substitute	(	const Array< Range > &	region,
		const Map< Var, PrimExpr > &	vmap
	)

Substitute the var specified by vmap.

Parameters

region	The object whose vars are to be substituted
vmap	The map of new values.

Returns: The result.

◆ Substitute() [4/5]

template<typename T >

auto tvm::tir::Substitute	(	T	input,
		const Map< Var, PrimExpr > &	value_map
	)

inline

Sugar for substitute via a given map.

Parameters

input	The input to be updated.
value_map	The map of new values.

Returns: The result.

Template Parameters

T	the input type, can be PrimExpr or Stmt.

◆ Substitute() [5/5]

template<typename T >

T tvm::tir::Substitute	(	T	input,
		const std::unordered_map< const VarNode *, PrimExpr > &	value_map
	)

inline

Sugar for substitute via a given map.

Parameters

input	The input to be updated.
value_map	The map of new values.

Returns: The result.

Template Parameters

T	the input type, can be PrimExpr or Stmt.

◆ TypeAnnotation()

PrimExpr tvm::tir::TypeAnnotation	(	DataType	dtype,
		Span	span = `Span()`
	)

Create a type annotation expression.

Parameters

dtype	The data type
span	The location of this object in the source code.

Returns: Expr a expression with dtype.

◆ UndefinedVars() [1/2]

Array<Var> tvm::tir::UndefinedVars	(	const Stmt &	stmt,
		const Array< Var > &	defs
	)

Find undefined vars in the statement.

Parameters

stmt	The function to be checked.
defs	The vars that is defined.

Returns: Array of undefined vars.

◆ UndefinedVars() [2/2]

Array<Var> tvm::tir::UndefinedVars ( const PrimExpr & expr )

Find undefined vars in the expression.

Parameters

expr	The expression to be checked.

Returns: Array of undefined vars.

◆ UsesVar() [1/2]

bool tvm::tir::UsesVar	(	const Stmt &	stmt,
		std::function< bool(const VarNode *)>	vset_contains
	)

Whether the given Stmt uses any var in the given variable set.

Parameters

stmt	The Stmt to be checked.
vset_contains	The check function to see if a var is in the variable set.

Returns: Whether stmt uses any var in the given variable set.

◆ UsesVar() [2/2]

bool tvm::tir::UsesVar	(	const PrimExpr &	expr,
		std::function< bool(const VarNode *)>	vset_contains
	)

Whether the given PrimExpr uses any var in the given variable set.

Parameters

expr	The PrimExpr to be checked.
vset_contains	The check function to see if var is in the variable set.

Returns: Whether expr uses any var in the given variable set.

◆ VerifyGPUCode()

bool tvm::tir::VerifyGPUCode	(	const PrimFunc &	func,
		Map< String, PrimExpr >	constraints
	)

Verify the correctness of a GPU code It will check the whether the amount of memory usage or the number of threads in a block exceeds the limit.

Parameters

func	The function to be checked
constraints	The dict to specify constraints to check. Possible keys are

"max_local_memory_per_block": Total amount of local memory per block (in bytes). "max_shared_memory_per_block": Total amount of shared memory per block (in bytes). "max_threads_per_block": Maximum number of threads per block. "max_thread_x": Maximum length of threadIdx.x. "max_thread_y": Maximum length of threadIdx.y. "max_thread_z": Maximum length of threadIdx.z.

If one key is missing in this argument, the pass won't check for that item.

Returns: valid Whether it is a valid GPU code

◆ VerifyMemory()

bool tvm::tir::VerifyMemory ( const PrimFunc & func )

Verify if memory accesses are legal for a specific target device type.

In the case that tgt is cuda, if not all workload is bound with threads, CPU code is generated that tries to access GPU memory, which is illegal. This pass performs verification for this case.

Parameters

func	The function to be verified.

Returns: Success of memory verification.

◆ VerifySSA()

bool tvm::tir::VerifySSA ( const PrimFunc & func )

Verifies whether the IR stmt or Expr is in SSA form. That is: each Var is defined and assigned once(in Let/For)

Parameters

func	The function to be verified.

Returns: Whether IR is in SSA form.

Note: All passes in TIR consume and produce SSA form.

◆ VisitPrimFuncs()

template<class FLambda >

void tvm::tir::VisitPrimFuncs	(	const IRModule &	mod,
		FLambda	fvisit
	)

inline

Visit the PrimFuncs in the IRModule.

Template Parameters

FLambda The type of the PrimFunc visitor

Parameters

mod	The IRModule to be visited
fvisit	The visitor to the PrimFuncs in the IRModule

Namespaces

Classes

Typedefs

Enumerations

Functions

Typedef Documentation

◆ ExprRV

◆ ExprRVNode

◆ FInstructionApply

◆ FInstructionAsPython

◆ FInstructionAttrsAsJSON

◆ FInstructionAttrsFromJSON

◆ FLegalize

◆ FloatImmNode

◆ FLowerIntrinsic

◆ FTraceDecisionProvider

◆ IntImmNode

◆ Region

◆ TCallEffectKind

◆ TGlobalSymbol

◆ TVectorizable

Enumeration Type Documentation

◆ BufferIndexType

◆ BufferType

◆ CallEffectKind

◆ DepKind

◆ ForKind

◆ IterVarType

◆ ScheduleDebugMask

◆ ScheduleErrorRenderLevel

Function Documentation

◆ as_const_int()

◆ as_unordered_map()

◆ CalculateExprComplexity()

◆ CalculateWorkspaceBytes()

◆ const_false()

◆ const_true()

◆ decl_buffer()

◆ DetectBufferAccessLCA()

◆ EstimateTIRFlops() [1/2]

◆ EstimateTIRFlops() [2/2]

◆ foldl()

◆ ForKind2String()

◆ GetBlockAccessRegion()

◆ GetBlockReadWriteRegion()

◆ IRTransform()

◆ is_const_int() [1/2]

◆ is_const_int() [2/2]

◆ is_const_number()

◆ is_const_power_of_two_integer()

◆ is_negative_const()

◆ is_no_op()

◆ is_one()

◆ is_positive_const()

◆ is_zero()

◆ IsPointerType()

◆ IterVarType2String()

◆ make_const()

◆ make_zero()

◆ MakeConstScalar()

◆ operator<<()

◆ PostOrderVisit()

◆ PreOrderVisit()

◆ RenewDefs()

◆ SideEffect()

◆ Specialize()

◆ Substitute() [1/5]

◆ Substitute() [2/5]

◆ Substitute() [3/5]

◆ Substitute() [4/5]

◆ Substitute() [5/5]

◆ TypeAnnotation()

◆ UndefinedVars() [1/2]

◆ UndefinedVars() [2/2]

◆ UsesVar() [1/2]

◆ UsesVar() [2/2]

◆ VerifyGPUCode()

◆ VerifyMemory()

◆ VerifySSA()

◆ VisitPrimFuncs()