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, *, ...)

Copy a prepacked NZ region from UB to L1 using explicit storage pitches.

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, *, src_pitch, dst_pitch, dtype)¶

Copy a prepacked NZ region from UB to L1 using explicit storage pitches.

rows and cols describe this transfer. Pitches are the distance between C0-group starts in 32-byte blocks, independent of the region size. Both pointers must include the region origin and be 32-byte aligned. This operation does not initialize gaps or convert element types.

Parameters:

dtype (str)

Return type:

None