tilelang.ascend.language.simd¶

T.simd.* - Raw CCE vector intrinsics for Ascend SIMD programming.

Function names match the underlying CCE intrinsics directly.

MODE_MERGING preserves inactive lanes of the mutable destination register. On Ascend 950, validated 8/16/32-bit operations use the CCE merging overloads; vdupv maps to CCE’s vector vdup overload. Scalar BF16 vdup retains software merging to work around CANN 9.2’s inactive-lane bug. Precision-specific SFU algorithms also retain their wrappers, including the default exact FP32 division and the ftz_false variants of vexp, vln, and vsqrt.

Classes¶

SimdPair

Wraps a two-result SIMD intrinsic and preserves each result dtype.

Functions¶

`pset`(elem_width[, dist])	Create a predicate mask: pset_bXX(dist).
`pge`(elem_width[, dist])	Create a predicate mask from pge_bXX(dist).
`update_mask`(value[, width])	Runtime tail predicate: lanes [0, value) active (b8/b16/b32).
`pand`(src0, src1, mask)
`por`(src0, src1, mask)
`pxor`(src0, src1, mask)
`pnot`(src, mask)
`psel`(src0, src1, mask)
`ppack`(src[, part])	Predicate pack 2:1 (zeroing): dst = ppack(src, LOWER/HIGHER).
`punpack`(src[, part])	Predicate unpack 1:2 (zeroing): dst = punpack(src, LOWER/HIGHER).
`pintlv`(src0, src1[, width])	Predicate interleave -> pair of predicates (b8/b16/b32).
`pdintlv`(src0, src1[, width])	Predicate deinterleave -> pair of predicates (b8/b16/b32).
`alloc_var`(dtype)	Allocate a single mutable SIMD register variable (return-value style).
`alloc_local`(shape, dtype)	Allocate an addressable array of mutable SIMD register variables.
`pld`(addr[, dist])	Load a predicate from UB. Express address offsets in `addr`.
`pst`(addr, src[, dist])	Store a predicate to UB. Express address offsets in `addr`.
`vld`(addr[, dist, post_inc])	Vector load. Returns a typed vector register.
`vld2`(addr[, dist, off])	Dual-dest vector load: `a, b = vld2(x_ub[i, col], dist="DINTLV_B8")`.
`vsts`(addr, src[, mask, dist, extent])	Vector store to `addr`.
`vsstb`(src, base, stride[, mask, update])	Scatter-store 32B blocks with an optional POST_UPDATE pointer.
`make_ubuf_ptr`(buf_access, dtype)	Allocate a mutable UB pointer for post-update `vld()` / `vsstb()` calls.
`vadd`(src0, src1[, mask, mode])
`vaddc`(src0, src1[, mask])	Add int32/uint32 vectors without carry-in and return `(carry, result)`.
`vsubc`(src0, src1[, mask])	Subtract int32/uint32 without carry-in and return `(carry, result)`.
`vaddcs`(src0, src1, carrysrcp[, mask])	Add int32/uint32 with carry-in predicate and return `(carry, result)`.
`vsubcs`(src0, src1, carrysrcp[, mask])	Subtract int32/uint32 with carry-in predicate and return `(carry, result)`.
`vmull`(src0, src1[, mask])	Widening 32x32->64 multiply returning `(lo, hi)` (int32/uint32).
`vsub`(src0, src1[, mask, mode])
`vmul`(src0, src1[, mask, mode])
`vmula`(dst, src0, src1[, mask, mode])	Fused multiply-add: dst = dst + src0 * src1.
`vmadd`(dst, src0, src1[, mask, mode])	Fused multiply-add: dst = dst * src0 + src1.
`vaxpy`(dst, src, scalar[, mask, mode])	Fused scalar multiply-add: dst = src * scalar + dst.
`dhistv2`(dst, src[, mask, bin])	Accumulate a frequency histogram of a uint8 vector into a uint16 vector.
`chistv2`(dst, src[, mask, bin])	Accumulate a cumulative histogram of a uint8 vector into a uint16 vector.
`vdiv`(src0, src1[, mask, mode, precision])	Divide vectors, optionally selecting the SFU implementation.
`vmax`(src0, src1[, mask, mode])
`vmin`(src0, src1[, mask, mode])
`vand`(src0, src1[, mask, mode])
`vor`(src0, src1[, mask, mode])
`vxor`(src0, src1[, mask, mode])
`vshl`(src0, src1[, mask, mode])
`vshr`(src0, src1[, mask, mode])
`vexp`(src[, mask, mode, precision])
`vln`(src[, mask, mode, precision])
`vsqrt`(src[, mask, mode, precision])
`vabs`(src[, mask, mode])
`vneg`(src[, mask, mode])
`vrelu`(src[, mask, mode])
`vlrelu`(src, alpha[, mask])	Leaky ReLU with scalar slope (f16/f32): dst = src >= 0 ? src : alpha * src.
`vprelu`(src0, src1[, mask])	Parametric ReLU with per-lane slope vector (f16/f32).
`vnot`(src[, mask, mode])
`vdup`(src, dtype_str[, mask, mode])	Broadcast scalar to all lanes: dst = vdup(scalar, dtype_str, mask).
`vdupv`(src[, mask, pos, mode])	Broadcast lane N of src vector to all lanes: dst = vdupv(src, mask, pos).
`vcpadd`(src[, mask, mode])	Pairwise adjacent-lane add.
`vcadd`(src[, mask, mode])	Pairwise add reduction: dst = vcadd(src, mask).
`vcmax`(src[, mask, mode])	Pairwise max reduction: dst = vcmax(src, mask).
`vcmin`(src[, mask, mode])	Pairwise min reduction: dst = vcmin(src, mask).
`vcgadd`(src[, mask, mode])	Grouped add reduction: dst = vcgadd(src, mask).
`vcgmax`(src[, mask, mode])	Grouped max reduction: dst = vcgmax(src, mask).
`vcgmin`(src[, mask, mode])	Grouped min reduction: dst = vcgmin(src, mask).
`vsqz`(src[, mask, mode])	Squeeze selected lanes toward the low lanes: dst = vsqz(src, mask).
`vusqz`(mask[, dtype])	Per-lane exclusive prefix count of mask (s8/s16/s32).
`vci`(index, dtype_str[, order])	Index ramp: dst = vci(index, dtype_str, order).
`vcmp`(src0, src1[, mask, op])	Elementwise compare -> vector_bool: dst = vcmp_<op>(src0, src1, mask).
`vcmps`(src, scalar[, mask, op])	Compare vector vs scalar -> vector_bool: dst = vcmps_<op>(src, scalar, mask).
`vintlv`(src0, src1)	Interleave two vector registers.
`vdintlv`(src0, src1)	De-interleave two vector registers.
`pair_get`(pair, index)	Extract element from a pair: reg = pair_get(pair, 0) or pair_get(pair, 1).
`vpack`(src[, part])	Pack wider lanes to narrower lanes: dst = vpack(src, LOWER/HIGHER).
`vunpack`(src[, part])	Widen half of src: u8->u16, s8->s16, u16->u32, s16->s32.
`vgatherb`(base, index[, mask])	Gather 32B blocks from base using vector_u32 block offsets. Returns a vector.
`vgather2`(base, index[, mask])	Gather elements from base using per-lane offsets.
`vscatter`(src, base, index[, mask])	Scatter-store: base[index[lane]] = src[lane].
`vexpdif`(src0, src1[, mask])	Fused exp-sub: dst = exp(src0 - src1).
`vabsdif`(src0, src1[, mask, mode])	Fused abs-sub: dst = vabsdif(src0, src1, mask, mode).
`vcvt`(src, target_dtype[, mask, round, sat, part, mode])	Vector type conversion between float32, float16, bfloat16, float8, float4, and integers.
`vsel`(src0, src1, mask)	Bitwise select: mask ? src0 : src1.
`vselr`(src, index)	Select lanes from src using per-lane indices.
`vmaxs`(src, scalar[, mask, mode])
`vmins`(src, scalar[, mask, mode])
`vmuls`(src, scalar[, mask, mode])
`vadds`(src, scalar[, mask, mode])
`vshls`(src, scalar[, mask, mode])
`vshrs`(src, scalar[, mask, mode])
`mem_bar`(mem_type)	Memory barrier: mem_bar(mem_type), e.g. mem_bar(VST_VLD).

Module Contents¶

class tilelang.ascend.language.simd.SimdPair(pair, dtype=None)¶

Wraps a two-result SIMD intrinsic and preserves each result dtype.

a, b = ... emits two pair_get calls against the pair. The pair- producing op (e.g. vintlv, vld2, or post-update vld) is bound once at its call site by the frontend, so both pair_get calls reference a single bound variable rather than inlining the pair expression twice.

dtype is a pair describing both result types, such as the (boolx256, int32x64) carry/result pair returned by vaddc. Omitting it defaults both result types to the dtype of the backing TIR expression.

property dtype¶

__getitem__(index)¶

__iter__()¶

tilelang.ascend.language.simd.pset(elem_width, dist='PAT_ALL')¶

Create a predicate mask: pset_bXX(dist).

Returns a vector_bool (boolx256). dist: “PAT_ALL”, “PAT_VL1”..”PAT_VL128”, “PAT_M3”, “PAT_M4”, “PAT_H”, “PAT_Q”, etc.

Parameters:

elem_width (int)
dist (str)

tilelang.ascend.language.simd.pge(elem_width, dist='PAT_ALL')¶

Create a predicate mask from pge_bXX(dist).

Parameters:

elem_width (int)
dist (str)

tilelang.ascend.language.simd.update_mask(value, width=32)¶: Runtime tail predicate: lanes [0, value) active (b8/b16/b32).

tilelang.ascend.language.simd.pand(src0, src1, mask)¶

tilelang.ascend.language.simd.por(src0, src1, mask)¶

tilelang.ascend.language.simd.pxor(src0, src1, mask)¶

tilelang.ascend.language.simd.pnot(src, mask)¶

tilelang.ascend.language.simd.psel(src0, src1, mask)¶

tilelang.ascend.language.simd.ppack(src, part=0)¶: Predicate pack 2:1 (zeroing): dst = ppack(src, LOWER/HIGHER).

tilelang.ascend.language.simd.punpack(src, part=0)¶: Predicate unpack 1:2 (zeroing): dst = punpack(src, LOWER/HIGHER).

tilelang.ascend.language.simd.pintlv(src0, src1, width=32)¶: Predicate interleave -> pair of predicates (b8/b16/b32).

tilelang.ascend.language.simd.pdintlv(src0, src1, width=32)¶: Predicate deinterleave -> pair of predicates (b8/b16/b32).

tilelang.ascend.language.simd.alloc_var(dtype)¶

Allocate a single mutable SIMD register variable (return-value style).

Uses local.var scope and behaves as one vector register value. For an addressable array of registers (v[i]), use alloc_local().

Parameters:: dtype (tilelang._typing.DType)

tilelang.ascend.language.simd.alloc_local(shape, dtype)¶

Allocate an addressable array of mutable SIMD register variables.

Uses local scope so each element is an individually addressable register, allowing indexed access like v[i]:

v = T.simd.alloc_local(4, "float32")
for i in T.Unroll(4, explicit=True):
    v[i] = vld(s_ub[i * VL])

Parameters:

shape (tilelang._typing.ShapeType)
dtype (tilelang._typing.DType)

tilelang.ascend.language.simd.pld(addr, dist='NORM')¶: Load a predicate from UB. Express address offsets in addr.

tilelang.ascend.language.simd.pst(addr, src, dist='NORM')¶: Store a predicate to UB. Express address offsets in addr.

tilelang.ascend.language.simd.vld(addr, dist='NORM', *, post_inc=None)¶

Vector load. Returns a typed vector register.

addr can be a BufferLoad auto-wrapped as tl.access_ptr. Express address offsets in addr.

With post_inc=step, load through a mutable make_ubuf_ptr() handle and return (vector, advanced_pointer). Assign the second result back to the handle. step is a signed int32 increment in elements of the dtype declared by make_ubuf_ptr(); the load uses the old address. The dtype must be 8/16/32-bit and match the distribution. None selects an ordinary load; zero still returns the pair without advancing the pointer:

src_ptr = T.simd.make_ubuf_ptr(T.access_ptr(src_ub[0], "r", extent=256), "uint16")
first, src_ptr = T.simd.vld(src_ptr, post_inc=128)
second, src_ptr = T.simd.vld(src_ptr, post_inc=128)

Keep the pointer within one SIMD VF and declare its complete accessed span in the initializer’s access_ptr, as in the example.

A BRC_B8/B16/B32 broadcast replicates one element of the width named by the suffix, widening the result past the source buffer’s element type when the two differ (BRC_B16 over a uint8 buffer broadcasts a 16-bit element, not a byte). Every other distribution takes its element width from the source buffer; their _B* suffixes describe the data being loaded and must agree with it.

tilelang.ascend.language.simd.vld2(addr, dist='DINTLV_B16', off=None)¶

Dual-dest vector load: a, b = vld2(x_ub[i, col], dist="DINTLV_B8").

Supported dists:

DINTLV_B8: load 512xu8/fp8 -> two 256-lane regs (even/odd bytes)
DINTLV_B16: load 256xbf16/u16 -> two 128-lane regs
DINTLV_B32: load 128xf32/u32 -> two 64-lane regs

The access_ptr footprint is always 2x the single-vector width. This is an opaque memory load: the pair is bound once at this program point so the two pair_get calls from the unpack share a single load rather than issuing two independent loads.

tilelang.ascend.language.simd.vsts(addr, src, mask=None, dist='NORM_B32', extent=None)¶

Vector store to addr.

addr can be a BufferLoad auto-wrapped as tl.access_ptr. Express address offsets in addr. Optional extent overrides the default access_ptr footprint (e.g. 8 for a dense PAT_VL8 NORM_B16 recip pack).

tilelang.ascend.language.simd.vsstb(src, base, stride, mask=None, update=False)¶

Scatter-store 32B blocks with an optional POST_UPDATE pointer.

Passing a regular buffer access performs a store and returns void. Set update=True with the mutable handle returned by make_ubuf_ptr() to enable POST_UPDATE and return the advanced pointer, which should be assigned back to the same handle:

dst_ptr = T.simd.make_ubuf_ptr(dst_ub[0], "bfloat16")
dst_ptr = T.simd.vsstb(src, dst_ptr, stride, mask, update=True)

tilelang.ascend.language.simd.make_ubuf_ptr(buf_access, dtype)¶

Allocate a mutable UB pointer for post-update vld() / vsstb() calls.

dtype declares the pointee element type used by vld() and checked against the source vector by vsstb(). It is stored in the enclosing block’s IR annotations; the mutable carrier remains a handle buffer.

The pointer is carried by a local.var handle buffer. Assigning the advanced handle returned by vld() or vsstb() writes it back into the same mutable carrier:

dst_ptr = T.simd.make_ubuf_ptr(dst_ub[0], "bfloat16")
dst_ptr = T.simd.vsstb(src, dst_ptr, stride, mask, update=True)

tilelang.ascend.language.simd.vadd(src0, src1, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vaddc(src0, src1, mask=None)¶

Add int32/uint32 vectors without carry-in and return (carry, result).

carry is a boolx256 predicate register and result has the same dtype as the inputs. The underlying Ascend vaddc instruction only supports full-register int32x64 and uint32x64 operands.

tilelang.ascend.language.simd.vsubc(src0, src1, mask=None)¶

Subtract int32/uint32 without carry-in and return (carry, result).

carry is 1 where the subtraction completes without borrow.

tilelang.ascend.language.simd.vaddcs(src0, src1, carrysrcp, mask=None)¶: Add int32/uint32 with carry-in predicate and return (carry, result).

tilelang.ascend.language.simd.vsubcs(src0, src1, carrysrcp, mask=None)¶: Subtract int32/uint32 with carry-in predicate and return (carry, result).

tilelang.ascend.language.simd.vmull(src0, src1, mask=None)¶: Widening 32x32->64 multiply returning (lo, hi) (int32/uint32).

tilelang.ascend.language.simd.vsub(src0, src1, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vmul(src0, src1, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vmula(dst, src0, src1, mask=None, mode='MODE_ZEROING')¶: Fused multiply-add: dst = dst + src0 * src1.

tilelang.ascend.language.simd.vmadd(dst, src0, src1, mask=None, mode='MODE_ZEROING')¶: Fused multiply-add: dst = dst * src0 + src1.

tilelang.ascend.language.simd.vaxpy(dst, src, scalar, mask=None, mode='MODE_ZEROING')¶: Fused scalar multiply-add: dst = src * scalar + dst.

tilelang.ascend.language.simd.dhistv2(dst, src, mask=None, bin=0)¶

Accumulate a frequency histogram of a uint8 vector into a uint16 vector.

bin=0 counts values in [0, 127] and bin=1 counts values in [128, 255]. The destination register is updated in place.

tilelang.ascend.language.simd.chistv2(dst, src, mask=None, bin=0)¶

Accumulate a cumulative histogram of a uint8 vector into a uint16 vector.

bin=0 returns cumulative counts through values [0, 127] and bin=1 returns cumulative counts through values [128, 255]. The destination register is updated in place.

tilelang.ascend.language.simd.vdiv(src0, src1, mask=None, mode='MODE_ZEROING', precision=None)¶

Divide vectors, optionally selecting the SFU implementation.

precision=None follows tl.enable_fast_math (precise fp32 division when fast math is off, hardware instruction otherwise). Non-fp32 division always uses the hardware instruction.

precision selects the implementation per op:

'ftz_true': bare hardware SFU (flush-to-zero semantics)
'exact' (alias 'vdiv_0ulp_ftz_true'): correctly-rounded fp32 division (CANN DivAlgo::PRECISION_0ULP_FTZ_TRUE / DivPrecisionImpl)

Requires float32 for precise paths.

tilelang.ascend.language.simd.vmax(src0, src1, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vmin(src0, src1, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vand(src0, src1, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vor(src0, src1, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vxor(src0, src1, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vshl(src0, src1, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vshr(src0, src1, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vexp(src, mask=None, mode='MODE_ZEROING', precision=None)¶

tilelang.ascend.language.simd.vln(src, mask=None, mode='MODE_ZEROING', precision=None)¶

tilelang.ascend.language.simd.vsqrt(src, mask=None, mode='MODE_ZEROING', precision=None)¶

tilelang.ascend.language.simd.vabs(src, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vneg(src, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vrelu(src, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vlrelu(src, alpha, mask=None)¶: Leaky ReLU with scalar slope (f16/f32): dst = src >= 0 ? src : alpha * src.

tilelang.ascend.language.simd.vprelu(src0, src1, mask=None)¶: Parametric ReLU with per-lane slope vector (f16/f32).

tilelang.ascend.language.simd.vnot(src, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vdup(src, dtype_str, mask=None, mode='MODE_ZEROING')¶

Broadcast scalar to all lanes: dst = vdup(scalar, dtype_str, mask).

dtype_str specifies the target vector element type (e.g. “float32”).

Parameters:: dtype_str (tilelang._typing.DType)

tilelang.ascend.language.simd.vdupv(src, mask=None, pos='POS_LOWEST', mode='MODE_ZEROING')¶

Broadcast lane N of src vector to all lanes: dst = vdupv(src, mask, pos).

pos: “POS_LOWEST” (lane 0) or “POS_HIGHEST” (lane N).

tilelang.ascend.language.simd.vcpadd(src, mask=None, mode='MODE_ZEROING')¶

Pairwise adjacent-lane add.

The sums of adjacent source lanes are packed into the low half of the result vector. Supported types: float32, float16.

tilelang.ascend.language.simd.vcadd(src, mask=None, mode='MODE_ZEROING')¶: Pairwise add reduction: dst = vcadd(src, mask).

tilelang.ascend.language.simd.vcmax(src, mask=None, mode='MODE_ZEROING')¶: Pairwise max reduction: dst = vcmax(src, mask).

tilelang.ascend.language.simd.vcmin(src, mask=None, mode='MODE_ZEROING')¶: Pairwise min reduction: dst = vcmin(src, mask).

tilelang.ascend.language.simd.vcgadd(src, mask=None, mode='MODE_ZEROING')¶: Grouped add reduction: dst = vcgadd(src, mask).

tilelang.ascend.language.simd.vcgmax(src, mask=None, mode='MODE_ZEROING')¶: Grouped max reduction: dst = vcgmax(src, mask).

tilelang.ascend.language.simd.vcgmin(src, mask=None, mode='MODE_ZEROING')¶: Grouped min reduction: dst = vcgmin(src, mask).

tilelang.ascend.language.simd.vsqz(src, mask=None, mode='MODE_STORED')¶: Squeeze selected lanes toward the low lanes: dst = vsqz(src, mask).

tilelang.ascend.language.simd.vusqz(mask, dtype='int32')¶: Per-lane exclusive prefix count of mask (s8/s16/s32).

tilelang.ascend.language.simd.vci(index, dtype_str, order='INC_ORDER')¶

Index ramp: dst = vci(index, dtype_str, order).

dst[lane] = index + lane (INC_ORDER) / index - lane (DEC_ORDER). dtype_str is the destination vector element type (e.g. “int32”, “float32”).

Parameters:: dtype_str (tilelang._typing.DType)

tilelang.ascend.language.simd.vcmp(src0, src1, mask=None, op='eq')¶

Elementwise compare -> vector_bool: dst = vcmp_<op>(src0, src1, mask).

op: eq/ne/gt/ge/lt/le.

tilelang.ascend.language.simd.vcmps(src, scalar, mask=None, op='lt')¶

Compare vector vs scalar -> vector_bool: dst = vcmps_<op>(src, scalar, mask).

op: eq/ne/gt/ge/lt/le.

tilelang.ascend.language.simd.vintlv(src0, src1)¶

Interleave two vector registers.

Unpack: a, b = vintlv(x, y)

The pair is bound once at the call site so the unpack’s two pair_get calls share a single permutation instruction instead of inlining the pair expression twice.

tilelang.ascend.language.simd.vdintlv(src0, src1)¶

De-interleave two vector registers.

Unpack: a, b = vdintlv(x, y)

The pair is bound once at the call site so the unpack’s two pair_get calls share a single permutation instruction instead of inlining the pair expression twice.

tilelang.ascend.language.simd.pair_get(pair, index)¶: Extract element from a pair: reg = pair_get(pair, 0) or pair_get(pair, 1).

tilelang.ascend.language.simd.vpack(src, part=0)¶

Pack wider lanes to narrower lanes: dst = vpack(src, LOWER/HIGHER).

Supports u32->u16 and u16->u8 (needed for dense UE8M0 scale packing).

tilelang.ascend.language.simd.vunpack(src, part=0)¶: Widen half of src: u8->u16, s8->s16, u16->u32, s16->s32.

tilelang.ascend.language.simd.vgatherb(base, index, mask=None)¶: Gather 32B blocks from base using vector_u32 block offsets. Returns a vector.

tilelang.ascend.language.simd.vgather2(base, index, mask=None)¶: Gather elements from base using per-lane offsets.

tilelang.ascend.language.simd.vscatter(src, base, index, mask=None)¶

Scatter-store: base[index[lane]] = src[lane].

Write-side counterpart to vgatherb/vgather2. index is a vector register of per-lane element offsets (uint32 for 32-bit elements, uint16 for 8/16-bit).

tilelang.ascend.language.simd.vexpdif(src0, src1, mask=None)¶

Fused exp-sub: dst = exp(src0 - src1).

The same-width form supports matching float32 vectors. Widening float16 inputs to float32 requires separate even/odd results and is not represented by this API.

tilelang.ascend.language.simd.vabsdif(src0, src1, mask=None, mode='MODE_ZEROING')¶

Fused abs-sub: dst = vabsdif(src0, src1, mask, mode).

Computes dst = abs(src0 - src1) in a single instruction. Supported types: float32, float16.

tilelang.ascend.language.simd.vcvt(src, target_dtype, mask=None, round='ROUND_R', sat=True, part=0, mode='MODE_ZEROING')¶

Vector type conversion between float32, float16, bfloat16, float8, float4, and integers.

Parameters:

src (VReg) – Source vector register holding the input element values.
target_dtype (T.dtype) – Destination element type, e.g. “float32”, “float16”, “bfloat16”, “float8_e4m3”, “float8_e5m2”, “float4_e2m1fn”, “int32”, “int16”, “uint16”, “int8”, “uint8”, “int64”.
mask (VReg (bool)) – Predicate mask; only lanes where mask[i] is true participate and are written. Defaults to an all-lanes mask matching the narrower (low-lane) element: the source width when widening, the target width when narrowing.
round (str) – IEEE-754 rounding mode. Valid values: - "ROUND_R" (default): Round to nearest, ties to even (banker’s rounding). - "ROUND_A": Round away from zero. - "ROUND_F": Round toward -inf (floor). - "ROUND_C": Round toward +inf (ceiling). - "ROUND_Z": Round toward zero (truncation). - "ROUND_O": Round to odd (only for f32->f16). - "ROUND_H": Round half away from zero (only for hif8 conversions). Not all modes are valid for every conversion pair; unused modes are ignored when absent.
sat (bool) – saturation mode for narrow/overflow-prone conversions. - True (default): Saturate to target range on overflow. - False: No sat; out-of-range values wrap. Appears in: all float->fp8, f32->f16/bf16, all float->int, integer narrowing. Absent from: widening conversions, f16->bf16, s16->f16, s32->f32 (no overflow possible).
part (int) – Sub-register half/quarter selector for widening/narrowing conversions. For even/odd 2-way splits: "PART_EVEN" (0), "PART_ODD" (1). For fp8/fp4/int4 4-way splits: "PART_P0" (0), "PART_P1" (1), "PART_P2" (2), "PART_P3" (3). Not needed for same-width conversions (e.g. f16<->bf16, s32->f32, f32->s32, f16->s16).
mode (str) – Write mode for masked-off lanes. - "MODE_ZEROING" (default): Inactive lanes are set to zero. - "MODE_MERGING": Inactive lanes preserve their prior value (only on mask-less int->int paths; 920R1 only for mask paths).

Returns:

VReg – Destination vector with the converted elements.
Supported conversion pairs (simplified)
—————————————-
* Float->Int (f32->s64/s32/s16, f16->s32/s16/s8/u8, bf16->s32.)
* Float->Float (f32->f16/bf16/fp8, f16->fp8/bf16/f32, bf16->fp8/f16/f32/fp4.)
* Int->Float (s16/s32/s64->f16/f32, s8/u8->f16.)
* Int->Int (Most s/u{8,16,32,64} widening/narrowing pairs.)

tilelang.ascend.language.simd.vsel(src0, src1, mask)¶: Bitwise select: mask ? src0 : src1.

tilelang.ascend.language.simd.vselr(src, index)¶: Select lanes from src using per-lane indices.

tilelang.ascend.language.simd.vmaxs(src, scalar, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vmins(src, scalar, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vmuls(src, scalar, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vadds(src, scalar, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vshls(src, scalar, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.vshrs(src, scalar, mask=None, mode='MODE_ZEROING')¶

tilelang.ascend.language.simd.mem_bar(mem_type)¶: Memory barrier: mem_bar(mem_type), e.g. mem_bar(VST_VLD).