tilelang.language.copy¶

The language interface for tl programs.

Functions¶

copy(src, dst[, coalesced_width, disable_tma, ...])

Copy data between memory regions.

c2d_im2col(img, col, nhw_step, c_step, kernel, stride, ...)

Perform im2col transformation for 2D convolution.

Module Contents¶

tilelang.language.copy.copy(src, dst, coalesced_width=None, disable_tma=False, eviction_policy=None)¶

Copy data between memory regions.

Parameters:

  • src (Union[tir.Buffer, tir.BufferLoad, tir.BufferRegion]) – Source memory region

  • dst (Union[tir.Buffer, tir.BufferLoad]) – Destination memory region

  • coalesced_width (Optional[int], optional) – Width for coalesced memory access. Defaults to None.

  • disable_tma (bool)

  • eviction_policy (Literal['evict_normal', 'evict_first', 'evict_last'] | None)

Raises:

TypeError – If copy extents cannot be deduced from arguments

Returns:

A handle to the copy operation

Return type:

tir.Call

Range handling notes: - Accepts Buffer/BufferRegion/BufferLoad on either side. Extents are

derived as follows: Buffer -> shape, BufferRegion -> [r.extent], BufferLoad -> extents from its inferred/encoded region.

  • If both src and dst are scalar BufferLoad without region extents, lowers to a direct store: dst[…] = src.

  • If one side is missing extents, it is treated as all-ones with the other side’s rank to enable broadcasting.

  • Extents are right-aligned and legalized via legalize_pairwise_extents: per tail-dimension, equal keeps as-is, a 1 broadcasts to the other, otherwise a conservative tir.max is used to remain safe for dynamic shapes.

  • The finalized extents are encoded with tl.region via to_buffer_region and passed through to the backend; low-level loop construction and any scope-specific decisions happen during lowering.

tilelang.language.copy.c2d_im2col(img, col, nhw_step, c_step, kernel, stride, dilation, pad, eviction_policy=None)¶

Perform im2col transformation for 2D convolution.

Parameters:

  • img (tir.Buffer) – Input image buffer

  • col (tir.Buffer) – Output column buffer

  • nhw_step (tir.PrimExpr) – Step size for batch and spatial dimensions

  • c_step (tir.PrimExpr) – Step size for channel dimension

  • kernel (int) – Kernel size

  • stride (int) – Stride of the convolution

  • dilation (int) – Dilation rate

  • pad (int) – Padding size

  • eviction_policy (Literal['evict_normal', 'evict_first', 'evict_last'] | None)

Returns:

A handle to the im2col operation

Return type:

tir.Call