tilelang.language.copy ====================== .. py:module:: tilelang.language.copy .. autoapi-nested-parse:: The language interface for tl programs. Functions --------- .. autoapisummary:: tilelang.language.copy.copy tilelang.language.copy.c2d_im2col Module Contents --------------- .. py:function:: copy(src, dst, coalesced_width = None, disable_tma = False, eviction_policy = None) Copy data between memory regions. :param src: Source memory region :type src: Union[tir.Buffer, tir.BufferLoad, tir.BufferRegion] :param dst: Destination memory region :type dst: Union[tir.Buffer, tir.BufferLoad] :param coalesced_width: Width for coalesced memory access. Defaults to None. :type coalesced_width: Optional[int], optional :raises TypeError: If copy extents cannot be deduced from arguments :returns: A handle to the copy operation :rtype: 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. .. py:function:: c2d_im2col(img, col, nhw_step, c_step, kernel, stride, dilation, pad, eviction_policy = None) Perform im2col transformation for 2D convolution. :param img: Input image buffer :type img: tir.Buffer :param col: Output column buffer :type col: tir.Buffer :param nhw_step: Step size for batch and spatial dimensions :type nhw_step: tir.PrimExpr :param c_step: Step size for channel dimension :type c_step: tir.PrimExpr :param kernel: Kernel size :type kernel: int :param stride: Stride of the convolution :type stride: int :param dilation: Dilation rate :type dilation: int :param pad: Padding size :type pad: int :returns: A handle to the im2col operation :rtype: tir.Call