tilelang.ascend.op.gemm.gemm_mad_blockscaled¶

Ascend block-scaled MAD GEMM lowering.

Attributes¶

GEMM_INST_MAD_BLOCK_SCALED

Classes¶

GemmMADBlockScaled

MXFP8 MAD with explicit SFA/SFB scale-factor operands.

Module Contents¶

tilelang.ascend.op.gemm.gemm_mad_blockscaled.GEMM_INST_MAD_BLOCK_SCALED = 'ascend.mad.blockscaled'¶
class tilelang.ascend.op.gemm.gemm_mad_blockscaled.GemmMADBlockScaled¶

Bases: tilelang.tileop.gemm_blockscaled.gemm_blockscaled_base.GemmBlockScaledMixin, tilelang.ascend.op.gemm.gemm_mad.GemmMAD

MXFP8 MAD with explicit SFA/SFB scale-factor operands.

L1 A/B inputs lower to tl.ascend_blockscaled_gemm_l1 with the scale pointers. L0 A/B inputs lower to tl.ascend_mad_mx: SFA/SFB are the MX slot handles of the data tiles (alloc_l0a_sf/alloc_l0b_sf), loaded by a preceding T.copy(sf_l1, view); the MAD reads the slots implied by its A/B data addresses, so the SF operands only contribute their read regions to scheduling.

property is_blockscaled: bool¶
Return type:

bool

infer_layout(target, thread_nums)¶
Parameters:

  • target (tvm.target.Target)

  • thread_nums (int)

lower(layout_map, target, thread_bounds, thread_index, mbar_phase_expr=None)¶
Parameters:

  • layout_map (dict)

  • target (tvm.target.Target)

  • thread_bounds (tvm.ir.Range)

  • thread_index (tvm.tirx.PrimExpr)

  • mbar_phase_expr (tvm.tirx.PrimExpr | None)