tilelang.ascend.language.vmi

TileLang VMI namespace for Ascend PTO kernels.

Attributes

vadd
vmul
vmax
vmin
vand
vor
vxor
vshl
vshr
vabs
vneg
vrelu
vexp
vln
vsqrt
vnot
vadds
vmuls
vmaxs
vmins
vshls
vshrs

Classes

VmiPair

Lazy pair wrapper for multi-result VMI calls.

Functions

alloc_local(shape, vreg_type)	Allocate a statically indexed array of VMI vector registers.
alloc_var(dtype, *, size)	Allocate a mutable VMI data register slot.
inside_vmi()	Return True when executing inside with T.SimdVF():.
require_vmi_scope([context])	Raise if a VMI op is used outside with T.SimdVF():.
vreg(lanes, dtype_like)	Return a logical VMI vector dtype descriptor.
mask(lanes)	Return a logical VMI predicate dtype descriptor.
vload(source[, offset, stride, block_stride, group, ...])	Load a VMI vector from either a_ub[0]-style buffer loads or ptr+offset form.
vstore(values, destination[, offset, mask, stride, ...])	Store one VMI vector to either a_ub[0]-style buffer loads or ptr+offset form.
vci(base, *, size[, order, loc, ip])
vdiv(lhs, rhs[, mask, precision, pmode, loc, ip])	Vector divide. precision='exact' expands CANN vdiv_0ulp_ftz_true.
vcmp(lhs, rhs, seed, cmp, *[, pmode, loc, ip])
vcmps(source, scalar, seed, cmp, *[, pmode, loc, ip])
vsel(mask, true_value, false_value, *[, pmode, loc, ip])
vselr(source, index, *[, loc, ip])
vbrc(value, *, size[, group, loc, ip])
vcadd(source, mask, *[, group, pmode, reassoc, loc, ip])
vcmax(source, mask, *[, group, pmode, loc, ip])
vcmin(source, mask, *[, group, pmode, loc, ip])
vcvt(source[, to_dtype, mask, rounding, saturate, ...])
vinterpret_cast(source[, to_dtype, loc, ip])
vexpdif(x, max_value, mask, *[, pmode, loc, ip])
vaxpy(x, acc, alpha, mask, *[, pmode, loc, ip])
vlrelu(x, slope, mask, *[, pmode, loc, ip])
vprelu(x, alpha, mask, *[, pmode, loc, ip])
vmull(a, b, mask, *[, pmode, loc, ip])
vmula(acc, lhs, rhs, mask, *[, pmode, loc, ip])
vdhist(acc, source, mask, *[, loc, ip])
vchist(acc, source, mask, *[, loc, ip])
vgather(source, offsets, mask, *[, pmode, loc, ip])
vscatter(value, destination, offsets, mask, *[, ...])
create_mask(active_lanes, *, size[, group, loc, ip])
vintlv(lhs, rhs, mask, *[, pmode, loc, ip])
vdintlv(lhs, rhs, mask, *[, pmode, loc, ip])
vunzip(source[, to_dtype, loc, ip])	Unzip every element into its low and high half-width halves.
vzip(low, high[, to_dtype, loc, ip])	Zip low/high half-width halves back (inverse of vunzip).

Module Contents

tilelang.ascend.language.vmi.alloc_local(shape, vreg_type)

Allocate a statically indexed array of VMI vector registers.

vreg_type is a numeric vector dtype produced by vreg(); bool/mask and handle element types are rejected. The resulting buffer may only be indexed by compile-time values, such as induction variables from T.unroll(..., explicit=True) on the PTO backend.

This API is valid only inside a T.SimdVF() scope: VMI allocations are SIMD vector registers, and ordinary T.alloc_local remains available outside that scope. A block-level allocation marker records that the array is an explicit VMI data allocation, so lowering and code generation do not infer that intent from the vector dtype of an ordinary SIMD allocation.

tilelang.ascend.language.vmi.alloc_var(dtype, *, size)

Allocate a mutable VMI data register slot.

The TIR buffer keeps the numeric element dtype (a single lane is the scalar dtype, more lanes are a vector dtype).

This API is valid only inside a T.SimdVF() scope: VMI allocations are SIMD vector registers, and the ordinary T.alloc_var remains available outside that scope. A block-level allocation marker records that the slot is an explicit VMI data register, so lowering and code generation do not have to infer that intent from the vector dtype or from a later store.

class tilelang.ascend.language.vmi.VmiPair(pair)

Lazy pair wrapper for multi-result VMI calls.

property dtype

__iter__()

__len__()

__getitem__(index)

tilelang.ascend.language.vmi.inside_vmi()

Return True when executing inside with T.SimdVF():.

Return type:

bool

tilelang.ascend.language.vmi.require_vmi_scope(context='T.vmi')

Raise if a VMI op is used outside with T.SimdVF():.

Parameters:

context (str)

Return type:

None

tilelang.ascend.language.vmi.vreg(lanes, dtype_like)

Return a logical VMI vector dtype descriptor.

Parameters:

lanes (int)

tilelang.ascend.language.vmi.mask(lanes)

Return a logical VMI predicate dtype descriptor.

Parameters:

lanes (int)

tilelang.ascend.language.vmi.vload(source, offset=None, *, size, stride=None, block_stride=None, group=None, loc=None, ip=None)

Load a VMI vector from either a_ub[0]-style buffer loads or ptr+offset form.

A VMI load always reads the full size lanes starting at the load address: the ISA intentionally keeps load instructions unmasked for performance, so masks apply to computation and stores, never to the load itself. Reading a full vector past a logical tail is still safe on device — the top 8K of UB is reserved as stack space, so the extra read-only lanes cannot fault; they carry undefined values that the compute and store masks must keep out of valid outputs.

tilelang.ascend.language.vmi.vstore(values, destination, offset=None, mask=None, *, stride=None, block_stride=None, group=None, loc=None, ip=None)

Store one VMI vector to either a_ub[0]-style buffer loads or ptr+offset form.

tilelang.ascend.language.vmi.vci(base, *, size, order=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vadd

tilelang.ascend.language.vmi.vmul

tilelang.ascend.language.vmi.vmax

tilelang.ascend.language.vmi.vmin

tilelang.ascend.language.vmi.vand

tilelang.ascend.language.vmi.vor

tilelang.ascend.language.vmi.vxor

tilelang.ascend.language.vmi.vshl

tilelang.ascend.language.vmi.vshr

tilelang.ascend.language.vmi.vabs

tilelang.ascend.language.vmi.vneg

tilelang.ascend.language.vmi.vrelu

tilelang.ascend.language.vmi.vexp

tilelang.ascend.language.vmi.vln

tilelang.ascend.language.vmi.vsqrt

tilelang.ascend.language.vmi.vnot

tilelang.ascend.language.vmi.vadds

tilelang.ascend.language.vmi.vmuls

tilelang.ascend.language.vmi.vmaxs

tilelang.ascend.language.vmi.vmins

tilelang.ascend.language.vmi.vshls

tilelang.ascend.language.vmi.vshrs

tilelang.ascend.language.vmi.vdiv(lhs, rhs, mask=None, *, precision=None, pmode=None, loc=None, ip=None)

Vector divide. precision='exact' expands CANN vdiv_0ulp_ftz_true.

tilelang.ascend.language.vmi.vcmp(lhs, rhs, seed, cmp, *, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vcmps(source, scalar, seed, cmp, *, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vsel(mask, true_value, false_value, *, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vselr(source, index, *, loc=None, ip=None)

tilelang.ascend.language.vmi.vbrc(value, *, size, group=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vcadd(source, mask, *, group=None, pmode=None, reassoc=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vcmax(source, mask, *, group=None, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vcmin(source, mask, *, group=None, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vcvt(source, to_dtype=None, mask=None, *, rounding=None, saturate=None, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vinterpret_cast(source, to_dtype=None, *, loc=None, ip=None)

tilelang.ascend.language.vmi.vexpdif(x, max_value, mask, *, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vaxpy(x, acc, alpha, mask, *, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vlrelu(x, slope, mask, *, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vprelu(x, alpha, mask, *, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vmull(a, b, mask, *, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vmula(acc, lhs, rhs, mask, *, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vdhist(acc, source, mask, *, loc=None, ip=None)

tilelang.ascend.language.vmi.vchist(acc, source, mask, *, loc=None, ip=None)

tilelang.ascend.language.vmi.vgather(source, offsets, mask, *, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vscatter(value, destination, offsets, mask, *, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.create_mask(active_lanes, *, size, group=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vintlv(lhs, rhs, mask, *, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vdintlv(lhs, rhs, mask, *, pmode=None, loc=None, ip=None)

tilelang.ascend.language.vmi.vunzip(source, to_dtype=None, *, loc=None, ip=None)

Unzip every element into its low and high half-width halves.

Returns (low, high) with the source lane count preserved and the element width halved; lane i of each result is the low/high half of source lane i. to_dtype selects the half element type and defaults to the unsigned half-width integer of the source element (uint16 for a 32-bit source). It is forwarded to the emitted PTO call, where PTODSL derives the half-width result element type from it and re-checks it against the source width, so it must agree with the result dtype this wrapper declares.

tilelang.ascend.language.vmi.vzip(low, high, to_dtype=None, *, loc=None, ip=None)

Zip low/high half-width halves back (inverse of vunzip).

to_dtype selects the wide element type and defaults to the unsigned double-width integer of the half element (uint32 for uint16 halves). As in vunzip it is forwarded to the emitted PTO call and must agree with the result dtype this wrapper declares.