Execution Scope IR

Execution Scope IR#

tvm.tirx.exec_scope#

Backend-defined execution scope identities.

class tvm.tirx.exec_scope.ExecScope(backend: str, name: str)

A node in one backend’s registered execution hierarchy.

tvm.tirx.exec_scope.register_exec_scope(backend: str, name: str, rank: int)

Register a backend’s finite scope hierarchy, from outer to inner ranks.

Backend load hooks register scopes before registering their layout axes.

Backend scopes#

Execution scope identities remain part of generic layout IR. Backends register their finite execution hierarchies and provide the user-facing constructors: T.cuda.ExecScope("warp") and T.trn.ExecScope("thread"). T.cuda.ExecScope("kernel") names the root used by grid axes such as bx in the kernel -> cta relationship. Every scope uniformly describes a hierarchy node; each backend tile instruction declares the scopes it supports. For example, CUDA TMA instructions accept thread and tcgen05.ld/st accept warpgroup. CUDA index bindings and predicate analysis live in tvm.backend.cuda.tile.exec_context.

CUDA execution hierarchy.

tvm.backend.cuda.exec_scope.ExecScope(name: str) → ExecScope

Construct a CUDA execution scope.

TRN execution hierarchy.

tvm.backend.trn.exec_scope.ExecScope(name: str) → ExecScope

Construct a TRN execution scope.

Contents