tvm_tirx.backend.cuda.layout.tmem_mma_operand_layout

tvm_tirx.backend.cuda.layout.tmem_mma_operand_layout(operand, shape, dtype, *, M, cta_group, ws=False, sparse=False, group=None)

Resolve the TMEM TileLayout for a tcgen05 MMA operand from operand role + instruction params.

operand: "A" (A-in-TMEM) or "D" (accumulator / C). M is the PTX tcgen05.mma instruction M (256/128/64), NOT the per-CTA buffer rows; per_cta_rows = M // cta_group. Physical 32-bit column count is left to the pool’s _resolve_cols (layout-aware).

Selects the documented PTX datapath from explicit instruction parameters.