tilelang.language.copy_op
=========================

.. py:module:: tilelang.language.copy_op

.. autoapi-nested-parse::

   Copy operations exposed on the TileLang language surface.



Functions
---------

.. autoapisummary::

   tilelang.language.copy_op.copy
   tilelang.language.copy_op.c2d_im2col


Module Contents
---------------

.. py:function:: copy(src, dst, *, coalesced_width = None, disable_tma = False, eviction_policy = None, annotations = None, loop_layout = 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, tir.BufferRegion]
   :param coalesced_width: Width for coalesced memory access. Defaults to None.
   :type coalesced_width: Optional[int], keyword-only
   :param disable_tma: Whether to disable TMA acceleration. Defaults to False.
   :type disable_tma: bool, keyword-only
   :param eviction_policy: Cache eviction policy. Defaults to None.
   :type eviction_policy: Optional[str], keyword-only
   :param annotations: Additional annotations dict. If provided,
                       coalesced_width, disable_tma, and eviction_policy can also be specified here.
                       Values in annotations take precedence over individual arguments.
   :type annotations: Optional[dict], keyword-only
   :param loop_layout: A parallel loop layout hint for the SIMT copy
                       (only valid for normal SIMT copy; incompatible with TMA/LDSM/STSM/TMem). When provided,
                       it is attached to the outermost parallel loop generated by this copy.
   :type loop_layout: Optional[Fragment], keyword-only

   :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