tilelang.ir¶

Classes¶

`Fill`
`AtomicAdd`
`Copy`
`Im2ColOp`
`GemmWarpPolicy`
`GemmSPWarpPolicy`
`FinalizeReducerOp`
`ParallelOp`
`ReduceOp`
`CumSumOp`
`CumMaxOp`
`ReduceType`
`Tcgen05Meta`	One tcgen05.ld/st data-movement shape: the per-warp single-issue
`Tcgen05CopyPlan`	The tiled copy of one TMEM tile built by tl.ExpandTcgen05Layout.

Functions¶

`make_span`(file, line)	Create a span covering a whole source line.
`set_stmt_span`(stmt, span)
`get_stmt_span`(stmt)
`set_buffer_span`(buffer, span)
`get_buffer_span`(buffer)
`set_prim_func_span`(func, span)
`stamp_stmt_spans`(stmt, span)	Fill in span on every Stmt in the subtree whose span is undefined.
`get_prim_func_span`(func)
`span_to_location`(span)	Convert a span to a (file, line) tuple, or None when undefined.
`span_coverage`(func)	Statistics on span injection coverage for a PrimFunc.

Module Contents¶

class tilelang.ir.Fill¶: Bases: tvm.ir.base.Node, tvm.runtime.Scriptable

class tilelang.ir.AtomicAdd¶: Bases: tvm.ir.base.Node, tvm.runtime.Scriptable

class tilelang.ir.Copy¶: Bases: tvm.ir.base.Node, tvm.runtime.Scriptable

class tilelang.ir.Im2ColOp¶: Bases: tvm.ir.base.Node, tvm.runtime.Scriptable

class tilelang.ir.GemmWarpPolicy¶

Bases: tvm.ir.base.Node, tvm.runtime.Scriptable

policy_type: int¶

m_warp: int¶

n_warp: int¶

compute_warp_partition(M, N, block_size, target, gemm_inst)¶

Parameters:

M (int)
N (int)
block_size (int)
target (tvm.target.Target)
gemm_inst (str)

class tilelang.ir.GemmSPWarpPolicy¶

Bases: tvm.ir.base.Node, tvm.runtime.Scriptable

policy_type: int¶

m_warp: int¶

n_warp: int¶

compute_warp_partition(M, N, block_size, target, gemm_inst)¶

Parameters:

M (int)
N (int)
block_size (int)
target (tvm.target.Target)
gemm_inst (str)

class tilelang.ir.FinalizeReducerOp¶: Bases: tvm.ir.base.Node, tvm.runtime.Scriptable

class tilelang.ir.ParallelOp¶: Bases: tvm.ir.base.Node, tvm.runtime.Scriptable

class tilelang.ir.ReduceOp¶: Bases: tvm.ir.base.Node, tvm.runtime.Scriptable

class tilelang.ir.CumSumOp¶: Bases: tvm.ir.base.Node, tvm.runtime.Scriptable

class tilelang.ir.CumMaxOp¶: Bases: tvm.ir.base.Node, tvm.runtime.Scriptable

class tilelang.ir.ReduceType¶: Bases: tvm.ir.base.Node, tvm.runtime.Scriptable

tilelang.ir.make_span(file, line)¶

Create a span covering a whole source line.

Parameters:

file (str)
line (int)

Return type:

tvm.ir.base.Span

tilelang.ir.set_stmt_span(stmt, span)¶

Parameters:: span (tvm.ir.base.Span)
Return type:: None

tilelang.ir.get_stmt_span(stmt)¶

Return type:: tvm.ir.base.Span | None

tilelang.ir.set_buffer_span(buffer, span)¶

Parameters:: span (tvm.ir.base.Span)
Return type:: None

tilelang.ir.get_buffer_span(buffer)¶

Return type:: tvm.ir.base.Span | None

tilelang.ir.set_prim_func_span(func, span)¶

Parameters:: span (tvm.ir.base.Span)
Return type:: None

tilelang.ir.stamp_stmt_spans(stmt, span)¶

Fill in span on every Stmt in the subtree whose span is undefined.

Nodes that already carry a span are left untouched. Intended for passes that rebuild a user statement into a fresh subtree: stamping the result with the original statement’s span keeps source-line information alive through lowering.

Parameters:: span (tvm.ir.base.Span | None)
Return type:: None

tilelang.ir.get_prim_func_span(func)¶

Return type:: tvm.ir.base.Span | None

tilelang.ir.span_to_location(span)¶

Convert a span to a (file, line) tuple, or None when undefined.

Parameters:: span (tvm.ir.base.Span | None)
Return type:: tuple[str, int] | None

tilelang.ir.span_coverage(func)¶

Statistics on span injection coverage for a PrimFunc.

Returns {“stmts”: [with_span, total], “buffers”: [with_span, total]}, where buffers cover both function parameters (buffer_map) and block- allocated buffers (SBlock.alloc_buffers, i.e. T.alloc_* sites). New nodes synthesized by passes legitimately have no span; this metric is meant for the freshly parsed IR (span injection validation).

Return type:: dict

class tilelang.ir.Tcgen05Meta¶

Bases: tvm.ir.base.Node, tvm.runtime.Scriptable

One tcgen05.ld/st data-movement shape: the per-warp single-issue CUTLASS TV atom plus the wrapper’s chaining limit.

intrinsics_name: str¶

max_chunks: int¶

class tilelang.ir.Tcgen05CopyPlan¶

Bases: tvm.ir.base.Node, tvm.runtime.Scriptable

The tiled copy of one TMEM tile built by tl.ExpandTcgen05Layout.

num_chunks_each_wg: int¶

num_issues: int¶

vals_per_issue: int¶