tilelang.contrib.ptodsl.gemm¶

Classes¶

PTOGemmL1Template

PTODSL helper for the dav-3510 Ascend L1 cube GEMM pipeline.

PTOBlockscaledGemmL1Template

PTODSL helper for an E4M3 or E2M1 blockscaled L1 cube GEMM tile.

Module Contents¶

class tilelang.contrib.ptodsl.gemm.PTOGemmL1Template(tile_m, tile_n, tile_k, base_k, sub_k_tiles, input_c0, sub_k_c0_blocks, a_l0_stage_elems, b_l0_stage_elems, input_pack_factor=1)¶

PTODSL helper for the dav-3510 Ascend L1 cube GEMM pipeline.

The constructor arguments are the static template parameters. Calling run_l1_tile emits the L1->L0 sub-K pipeline for one already-loaded L1 tile.

dav-3510 requires the TRANS_B=true path. The right operand must be staged in L1 as logical W[N, K]. If you write NN matmul with W[K, N], then during GM->L1 you must use DN2NZ to transpose B into this layout before calling the helper.

Parameters:

  • tile_m (int)

  • tile_n (int)

  • tile_k (int)

  • base_k (int)

  • sub_k_tiles (int)

  • input_c0 (int)

  • sub_k_c0_blocks (int)

  • a_l0_stage_elems (int)

  • b_l0_stage_elems (int)

  • input_pack_factor (int)

tile_m¶
tile_n¶
tile_k¶
base_k¶
sub_k_tiles¶
input_c0¶
sub_k_c0_blocks¶
a_l0_stage_elems¶
b_l0_stage_elems¶
input_pack_factor = 1¶
sub_k_storage_cols¶
run_l1_tile(a_mat, b_mat, a_l0_0, b_l0_0, acc, *, clear_accum, unit_flag_ctrl=0, tf32_mode=None)¶

Emit GEMM for one L1 A/B tile into acc.

class tilelang.contrib.ptodsl.gemm.PTOBlockscaledGemmL1Template(tile_m, tile_n, tile_k, base_k, sub_k_tiles, input_c0, sub_k_c0_blocks, a_l0_stage_elems, b_l0_stage_elems, sf_nz_stride, input_pack_factor=1)¶

Bases: PTOGemmL1Template

PTODSL helper for an E4M3 or E2M1 blockscaled L1 cube GEMM tile.

Matrix data uses the regular L1-to-L0 loads inherited from PTOGemmL1Template. Pair-packed uint16 scale storage is viewed as E8M0 only at the MX load boundary, where the caller supplies an pto.ptr(pto.f8e8m0, "mat") pointer to the same L1 allocation. FP4 uses input_pack_factor=2 so data-load start columns and L0 stage offsets remain in packed-storage units.

Parameters:

  • tile_m (int)

  • tile_n (int)

  • tile_k (int)

  • base_k (int)

  • sub_k_tiles (int)

  • input_c0 (int)

  • sub_k_c0_blocks (int)

  • a_l0_stage_elems (int)

  • b_l0_stage_elems (int)

  • sf_nz_stride (int)

  • input_pack_factor (int)

sf_nz_stride¶
sf_pairs_per_inner¶
run_l1_tile(a_mat, b_mat, sfa_e8m0_mat, sfb_e8m0_mat, a_l0_0, b_l0_0, acc, *, sf_k_offset, clear_accum, unit_flag_ctrl=0)¶

Emit one blockscaled L1 tile with E8M0 scale staging.