tvm_tirx.backend.cuda.script.tile.apply

Contents

tvm_tirx.backend.cuda.script.tile.apply#

tvm_tirx.backend.cuda.script.tile.apply(instruction, *operands, scope='thread', **parameters)#

Repeat the explicitly named PTX instruction over its canonical operands.

Operand axes are group, repeat, then instruction-specific atom coordinates. Group axes bind execution IDs; repeat axes generate unrolled loops. Each repeat emits the selected native instruction. The caller supplies any synchronization and asynchronous completion operations.

Parameters:

  • instruction (str or PTX instruction) – The exact PTX instruction, including its type and other qualifiers. TVMScript accepts, for example, T.ptx["ld.global.v4.b32"].

  • *operands (TensorRegion or Expr) – Operands in the instruction family’s documented order. A tensor denotes its full region. Scalar expressions are accepted where that family allows them; each tensor view must satisfy its storage and ownership requirements.

  • scope (str, optional) – Execution scope: "thread" (default), "warp", "warpgroup", "cta", or "cluster". It determines the leading group coordinates and does not launch threads or perform thread election.

  • **parameters – Family-specific static parameters and runtime operands. See the CUDA tile instruction reference for accepted keywords, defaults, required operands, and layout constraints. The generic IR construction keywords attrs, ty, ty_args, and loc are also accepted.

Returns:

call (tvm_tirx.ir.Call) – The tile operation as an IR call with a void result.

See also

tvm_tirx.backend.cuda.script.tile.ewise

Notes

The Tile Operands and Execution guide explains the coordinate conventions with a complete load, compute, and store example.

Contents