tilelang.ascend.language.kernel¶

Ascend NPU dialect of T.Kernel.

Ascend owns its launch surface. The NPU launch is a 1-D grid of AI cores with no SIMT thread domain at kernel scope, so this dialect’s Kernel declares prelude and nothing else: passing threads= or cluster_dims= is rejected by Python itself rather than by a runtime probe inside the shared launch path. Thread domains are declared explicitly inside the kernel body by T.SimtVF(threads=...) (real threadIdx scopes) or T.SimdVF() (register level, no threads); both emit their own thread scopes below the launch nest, so the kernel-level tx/ty/tz placeholders are dropped by the Ascend pipeline.

Classes¶

SimtVFContext

Thread vars and extents of the innermost nested thread scope.

Functions¶

`push_simtvf_context`(ctx)	Enter a nested thread scope, making its thread vars the current ones.
`pop_simtvf_context`()	Leave the innermost nested thread scope.
`get_thread_binding`([dim])	Returns the thread binding for the given dimension.
`get_thread_bindings`()	Returns all three thread bindings.
`get_thread_extent`([dim])	Returns the thread extent for the given dimension.
`get_thread_extents`()	Returns all three thread extents.
`Kernel`(*blocks[, prelude])	Construct a kernel launch frame for Ascend: a 1-D grid of AI cores.
`MixedKernel`(*blocks[, sids, prelude])	Construct an Ascend NPU Mixed-kernel launch frame with sub-kernel ID binding.

Module Contents¶

class tilelang.ascend.language.kernel.SimtVFContext(thread_vars, thread_extents)¶

Thread vars and extents of the innermost nested thread scope.

__slots__ = ('thread_vars', 'thread_extents')¶

thread_vars¶

thread_extents¶

tilelang.ascend.language.kernel.push_simtvf_context(ctx)¶

Enter a nested thread scope, making its thread vars the current ones.

Parameters:: ctx (SimtVFContext)

tilelang.ascend.language.kernel.pop_simtvf_context()¶: Leave the innermost nested thread scope.

tilelang.ascend.language.kernel.get_thread_binding(dim=0)¶

Returns the thread binding for the given dimension.

Parameters:: dim (int)
Return type:: tvm.tirx.Var

tilelang.ascend.language.kernel.get_thread_bindings()¶

Returns all three thread bindings.

Return type:: list[tvm.tirx.Var]

tilelang.ascend.language.kernel.get_thread_extent(dim=0)¶

Returns the thread extent for the given dimension.

Parameters:: dim (int)
Return type:: int

tilelang.ascend.language.kernel.get_thread_extents()¶

Returns all three thread extents.

Return type:: list[int]

tilelang.ascend.language.kernel.Kernel(*blocks, prelude=None)¶

Construct a kernel launch frame for Ascend: a 1-D grid of AI cores.

The grid becomes the NPU core index (blockIdx.x). There is no SIMT thread domain at this scope, so this dialect has no threads parameter: with T.Kernel(N) as bx yields one program index, and bx is iterable as (bx,). Use T.SimtVF(threads=...) inside the body to run thread-parallel code, or T.MixedKernel for an AIC+AIV mixed kernel.

Parameters:

*blocks (int | PrimExpr) – Extent of the 1-D core grid. Exactly one dimension is allowed; a multi-dimensional launch is rejected here rather than silently flattened downstream.
prelude (str, optional) – AscendC source injected before the generated kernel, e.g. #include lines or helper functions.

Return type:

tilelang.language.kernel.KernelLaunchFrame

Examples

with T.Kernel(NUM_CORES) as bx:
    with T.SimtVF(threads=128):
        for i in T.Parallel(128):
            out[bx * 128 + i] = x[bx * 128 + i] * 2.0

tilelang.ascend.language.kernel.MixedKernel(*blocks, sids=2, prelude=None)¶

Construct an Ascend NPU Mixed-kernel launch frame with sub-kernel ID binding.

Parameters:

blocks (int) – Number of blocks in the 1-D grid (blockIdx.x extent).
sids (int) – Number of active AIV sub-cores (1 or 2). Default is 2. On dav-3510, mixed kernels always launch the physical __mix__(1, 2) group; sids=1 restricts the vector body to sub-core 0. Binds as sid via asc_get_sub_block_id() in generated code.
prelude (str, optional) – Import C code injected before the generated kernel.

Returns:

res – The resulting frame providing (bx, sid) variable bindings.

Return type:

KernelLaunchFrame

Examples

with T.MixedKernel(NUM_BLOCKS, sids=2) as (bx, sid):
    with T.Cube():
        # AIC code using bx
        ...
    with T.Vector():
        # AIV code using both bx and sid
        ...