tilelang.cuda.language.copy_op¶

CUDA dialect of the copy operators: the common ops plus CUDA copy hints.

Functions¶

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

Copy data between memory regions, with CUDA lowering hints.

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

Perform im2col transformation for 2D convolution, with CUDA hints.

Module Contents¶

tilelang.cuda.language.copy_op.copy(src, dst, *, coalesced_width=None, disable_tma=False, eviction_policy=None, prefer_instruction=None, annotations=None, loop_layout=None)¶

Copy data between memory regions, with CUDA lowering hints.

Same semantics as the common tilelang.language.copy_op.copy(); the extra keywords steer how the CUDA backend lowers the copy. They are performance hints recorded on the tile op: compiling the same kernel for a target that has no use for them leaves the result unchanged.

Parameters:

  • src (tilelang._typing.BufferLikeType) – Source memory region (Buffer, BufferLoad or BufferRegion).

  • dst (tilelang._typing.BufferLikeType) – Destination memory region.

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

  • disable_tma (bool, keyword-only) – Never lower this copy through TMA even when the shape and scopes qualify. Defaults to False.

  • eviction_policy (Optional[str], keyword-only) – L2 cache eviction priority for the generated load/store or TMA instruction, one of "evict_normal", "evict_first", "evict_last".

  • prefer_instruction (Optional[str], keyword-only) – Preferred lowering instruction category: "tma", "cp_async" or "sync". For "tma", T.copy keeps synchronous copy semantics; global -> shared copies lower through TMA with an automatically allocated barrier and wait when constraints are satisfied.

  • annotations (Optional[dict], keyword-only) – Additional annotations dict; values in it take precedence over the individual keywords.

  • loop_layout (Optional[Fragment], keyword-only) – Parallel loop layout hint for the SIMT copy path.

Returns:

A handle to the copy operation.

Return type:

tirx.Call

tilelang.cuda.language.copy_op.im2col(img, col, nhw_step, c_step, kernel, stride, dilation, pad, eviction_policy=None, annotations=None)¶

Perform im2col transformation for 2D convolution, with CUDA hints.

Same semantics as the common tilelang.language.copy_op.im2col(); eviction_policy is the L2 cache hint consumed by the CUDA TMA im2col lowering (ignored by the generic SIMT fallback other targets use).

Parameters:

  • img (tilelang._typing.BufferLikeType)

  • col (tilelang._typing.BufferLikeType)

  • nhw_step (tvm.tirx.PrimExpr)

  • c_step (tvm.tirx.PrimExpr)

  • kernel (int)

  • stride (int)

  • dilation (int)

  • pad (int)

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

  • annotations (dict | None)

Return type:

tvm.tirx.PrimExpr