tilelang.layout.gemm_sp¶

Wrapping Layouts.

Functions¶

`decompose_col_major`(index_1d, basis)
`make_cutlass_metadata_layout_sm90`(buffer, mma_dtype, ...)	Make a layout of metadata that is compatible with cutlass sm90 compression kernel. Note that layout atom is the same for smem and gmem.
`make_cutlass_metadata_layout_sm8x`(buffer, mma_dtype)	Make a layout of metadata that is compatible with cutlass sm8x compression kernel. Note that layout atom is the same for smem and gmem.
`make_cutlass_metadata_layout`(buffer[, mma_dtype, arch])

tilelang.layout.gemm_sp.decompose_col_major(index_1d, basis)¶

Parameters:

Return type:

list[int]

tilelang.layout.gemm_sp.make_cutlass_metadata_layout_sm90(buffer, mma_dtype, block_k)¶

Make a layout of metadata that is compatible with cutlass sm90 compression kernel. Note that layout atom is the same for smem and gmem.

Parameters:

buffer (tvm.tir.Buffer) – metadata buffer shape, for sm90 it should be a 8-bit type
mma_dtype (str) – dtype of mma operand A, different dtypes result in different layout atom
block_k (int) – tiling size along K dim, different block_ks results in different layout atom.

tilelang.layout.gemm_sp.make_cutlass_metadata_layout_sm8x(buffer, mma_dtype)¶

Make a layout of metadata that is compatible with cutlass sm8x compression kernel. Note that layout atom is the same for smem and gmem.: ref: https://github.com/pytorch/pytorch/blob/d0c24b392cbb7b213d22e42c52c6c2d1ac2da1bd/torch/sparse/_semi_structured_conversions.py#L5

Parameters:

buffer (tvm.tir.Buffer) – metadata buffer shape, for sm80 it should be a 16bit type
mma_dtype (str)

tilelang.layout.gemm_sp.make_cutlass_metadata_layout(buffer, mma_dtype='float16', arch=None, **extra_args)¶

Parameters: