tilelang.ascend.language.dma

Ascend MTE DMA copy primitives and their pad-value / ND->NZ helpers.

Functions

ascend_set_copy_pad_value(value[, dtype])	Set the padding fill value for subsequent Ascend padded MTE copies.
ascend_nd2nz_post_copy(dst, src, rows, cols, ...)	Ascend ND->NZ post-copy from UB scratch to L1 with NZ pointer correction.

Module Contents

tilelang.ascend.language.dma.ascend_set_copy_pad_value(value, dtype=None)

Set the padding fill value for subsequent Ascend padded MTE copies.

Generates: asc_set_copy_pad_val(value_bits);

This call moves no data; it configures the stateful hardware pad value consumed by a following padded GM -> UB/L1 copy. Set it immediately before each padded copy rather than relying on a single set persisting across unrelated copies.

Parameters:

• value (int, float, or PrimExpr) – Padding fill value.

• dtype (str, optional) – Copy element dtype (one of int8/uint8/int16/uint16/float16/bfloat16/ int32/uint32/float32). Required when value is a Python literal; for a typed PrimExpr it defaults to value.dtype.

Return type:

None

Example

>>> T.ascend_set_copy_pad_value(0.0, dtype="float32")

tilelang.ascend.language.dma.ascend_nd2nz_post_copy(dst, src, rows, cols, full_rows, dst_dtype)

Ascend ND->NZ post-copy from UB scratch to L1 with NZ pointer correction.

Parameters:

dst_dtype (str)

Return type:

None