tilelang.ascend.language.frame

Execution-region frames for Ascend NPU.

Classes

CubeFrame	Frame for Cube (AIC) execution region.
VectorFrame	Frame for Vector (AIV) execution region.
SimtVFFrame	Frame for SimtVF (SIMT Vector Fixed) region.
SimdVFFrame	Frame for SimdVF (SIMD Vector Fragment) region.

Functions

Cube()	Construct a Cube (AIC) execution region frame.
Vector([vector])	Construct a Vector (AIV) execution region frame.
SimtVF([threads, latency])	Construct a dedicated SIMT vector-fragment region frame.
inside_simdvf()	Return whether the current builder is inside a SimdVF region.
SimdVF([latency])	Construct a SimdVF region for Ascend NPU MicroAPI operations.

Module Contents

class tilelang.ascend.language.frame.CubeFrame

Bases: tvm.tirx.script.builder.frame.TIRFrame

Frame for Cube (AIC) execution region.

tilelang.ascend.language.frame.Cube()

Construct a Cube (AIC) execution region frame.

Return type:

CubeFrame

class tilelang.ascend.language.frame.VectorFrame

Bases: tvm.tirx.script.builder.frame.TIRFrame

Frame for Vector (AIV) execution region.

In Mix mode, provides a subblock ID variable via as binding:

with T.Vector(vector=2) as sid:
    # sid = asc_get_sub_block_id()
    ...

__enter__()

tilelang.ascend.language.frame.Vector(vector=2)

Construct a Vector (AIV) execution region frame.

Parameters:

vector (int) – Number of AIV cores (1 or 2). Default 2.

Returns:

res – The VectorFrame used to denote an AIV execution region.

Return type:

VectorFrame

class tilelang.ascend.language.frame.SimtVFFrame

Bases: tvm.tirx.script.builder.frame.TIRFrame

Frame for SimtVF (SIMT Vector Fixed) region.

The C++ frame generates the SIMT_VF block and its thread-extent bindings. The Python frame additionally maintains the active SIMT context.

__enter__()

__exit__(ptype, value, trace)

tilelang.ascend.language.frame.SimtVF(threads=128, latency=0)

Construct a dedicated SIMT vector-fragment region frame.

Parameters:

• threads (int | list[int] | tuple) – Thread count for the SimtVF frame. A sequence may contain up to three dimensions; omitted dimensions default to one.

• latency (int) – Measured latency in cycles. Zero uses the scheduler’s latency model.

Returns:

res – The SimtVFFrame used to denote a SIMT vector-fragment region.

Return type:

SimtVFFrame

class tilelang.ascend.language.frame.SimdVFFrame

Bases: tvm.tirx.script.builder.frame.TIRFrame

Frame for SimdVF (SIMD Vector Fragment) region.

This is a register-level SIMD region without thread dimensions. The C++ frame generates the SIMD_VF block and its scope marker.

__enter__()

__exit__(ptype, value, trace)

tilelang.ascend.language.frame.inside_simdvf()

Return whether the current builder is inside a SimdVF region.

Return type:

bool

tilelang.ascend.language.frame.SimdVF(latency=0)

Construct a SimdVF region for Ascend NPU MicroAPI operations.

Parameters:

latency (int) – Measured latency in cycles. Zero uses the scheduler’s latency model.

Returns:

res – The SimdVFFrame used to denote a register-level SIMD region.

Return type:

SimdVFFrame