Thread Scope Enum Tags

tvm_tirx.tirx.thread_scope

Canonical enum tags for backend-defined thread hierarchies.

class tvm_tirx.tirx.thread_scope.ThreadScope(*args: Any, **kwargs: Any)

A registered thread scope, queried by its qualified backend.scope name.

Obtain a tag with ThreadScope.get("cuda.warp"), for example; unknown names raise an error. backend identifies its owner and rank orders the hierarchy, with smaller ranks denoting wider scopes. The registration index does not determine the hierarchy.

classmethod get(name: str) → ThreadScope

Look up a registered tag without creating unknown scopes.

property name: str

Canonical qualified name, such as cuda.warp.

property backend: str

Backend owning this tag.

property local_name: str

Backend-local name, such as warp.

property rank: int

Explicit hierarchy rank; smaller ranks denote wider scopes.

Backend scopes

Thread scopes are canonical ffi.Enum identities shared by layout IR, tile Call attributes, and dispatch contexts. Query a registered tag with T.ThreadScope("cuda.warp"). The name property is the qualified enum name; local_name is "warp". backend and rank are enum attributes, and registration order has no hierarchy semantics.

Backends register their finite thread hierarchies natively using EnumDef and explicit backend and rank attributes. Their user-facing lookup helpers are: T.cuda.ThreadScope("warp") and T.trn.ThreadScope("thread"). T.cuda.ThreadScope("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. Same-backend hierarchy comparisons use explicit ranks; cross-backend comparisons are rejected. For example, CUDA TMA instructions accept thread and tcgen05.ld/st accept warpgroup. CUDA index bindings and predicate analysis live in tvm_tirx.backend.cuda.tile.exec_context.

Instruction APIs continue accepting scope="warp" and also accept qualified names or ThreadScope values. Construction resolves the value before creating the Call; Call.attrs.scope always contains a ThreadScope enum. Omitting scope uses the instruction backend’s thread tag. Direct attribute construction must provide the tag or its qualified name; shared attribute schemas have no backend-independent default.

CUDA thread-scope lookup.

tvm_tirx.backend.cuda.thread_scope.ThreadScope(name: str) → ThreadScope

Look up a CUDA thread scope by its backend-local name.

TRN thread-scope lookup.

tvm_tirx.backend.trn.thread_scope.ThreadScope(name: str) → ThreadScope

Look up a TRN thread scope by its backend-local name.