tilelang.cuda.intrinsics.sparse_layout¶

Layout metadata for 2:4 semi-structured sparsity.

Shared contract between the example sparse compressor (examples.gemm_sp.sparse_utils) and the CUDA mma_sp code generator. Depends only on the leaf dtypes module so it can be pulled into the codegen import path without touching the language facade.

Attributes¶

GROUP_CONFIG

Functions¶

get_e_factor(a_dtype, meta_dtype)

Return how many a_dtype elements are indexed by one meta_dtype element.

get_e_replicate_factor(a_dtype)

Return how many consecutive threads share the same logical metadata value.

Module Contents¶

tilelang.cuda.intrinsics.sparse_layout.GROUP_CONFIG: dict[tilelang.language.dtypes.dtype, tuple[int, int]]¶
tilelang.cuda.intrinsics.sparse_layout.get_e_factor(a_dtype, meta_dtype)¶

Return how many a_dtype elements are indexed by one meta_dtype element.

Parameters:
Return type:

int

tilelang.cuda.intrinsics.sparse_layout.get_e_replicate_factor(a_dtype)¶

Return how many consecutive threads share the same logical metadata value.

Parameters:

a_dtype (tilelang.language.dtypes.dtype)

Return type:

int