tilelang.ir

Classes

Fill

AtomicAdd

Copy

Im2ColOp

GemmWarpPolicy

GemmSPWarpPolicy

FinalizeReducerOp

ParallelOp

ReduceOp

CumSumOp

CumMaxOp

RegionOp

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)

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.RegionOp

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.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