[Backend] Refactor gemm_sp

benchmark/matmul/benchmark_matmul_sp.py

@@ -4,21 +4,14 @@
 import torch
 from triton.testing import do_bench
 
-import tilelang
 import tilelang.language as T
 from tilelang.autotuner import autotune
 from tilelang import jit
-from tilelang.contrib import nvcc
-from tilelang.layout import make_cutlass_metadata_layout
+from tilelang.utils.sparse import get_e_factor
 
-# Configure logger
 logger = logging.getLogger(__name__)
 logger.setLevel(logging.DEBUG)
 
-arch = nvcc.get_target_compute_version()
-
-ARCH_INFO = {"8.0": (16, "int16"), "8.9": (16, "int16"), "9.0": (8, "uint8")}
-
 
 def ref_program(A, B):
     """
@@ -88,7 +81,7 @@ def get_configs(M, N, K):
     return configs
 
 
-def matmul_sp(M, N, K, in_dtype, accum_dtype):
+def matmul_sp(M, N, K, in_dtype, accum_dtype, e_dtype):
     """
     Create an autotuned matrix multiplication kernel for matrices of shape:
       - A: (M, K)
@@ -164,8 +157,8 @@ def kernel(
             A TVM Tensor Language function (T.prim_func) that computes matmul.
         """
         # Use half-precision for input data to reduce memory bandwidth,
-        # accumulate in float for better numerical accuracy
-        e_factor, e_dtype = ARCH_INFO[arch]
+        # accumulate in float for better numerical accurac
+        e_factor = get_e_factor(in_dtype, e_dtype)
 
         @T.prim_func
         def main(
@@ -200,15 +193,8 @@ def main(
 
                 # Clear out the accumulation buffer
                 T.clear(C_local)
-                T.disable_warp_group_reg_alloc()
 
                 T.use_swizzle(panel_size=10, enable=enable_rasterization)
-                T.annotate_layout(
-                    {
-                        E: make_cutlass_metadata_layout(E, mma_dtype=in_dtype, block_k=block_K),
-                        E_shared: make_cutlass_metadata_layout(E_shared, mma_dtype=in_dtype, block_k=block_K),
-                    }
-                )
                 # Loop over sub-blocks in K dimension, pipelined by num_stages
                 for k in T.Pipelined(T.ceildiv(K, block_K), num_stages=num_stages):
                     # Load a sub-block of A from global memory into A_shared
@@ -219,11 +205,13 @@ def main(
                     T.copy(B[k * block_K, bx * block_N], B_shared)
                     # Perform a partial matrix multiplication:
                     #   C_local += A_shared @ B_shared
-                    T.gemm_sp_v2(
+                    T.gemm_sp(
                         A_shared,
                         E_shared,
                         B_shared,
                         C_local,
+                        transpose_A=False,
+                        transpose_E=False,
                         transpose_B=False,
                         policy=policy,
                     )
@@ -242,8 +230,8 @@ def main(
     parser.add_argument("--m", type=int, default=16384, help="Matrix dimension M")
     parser.add_argument("--n", type=int, default=16384, help="Matrix dimension N")
     parser.add_argument("--k", type=int, default=16384, help="Matrix dimension K")
-    parser.add_argument("--disable_cache", action="store_true")
     parser.add_argument("--accum_dtype", type=str, default="float", choices=["float", "float16"], help="Accumulation datatype")
+    parser.add_argument("--e_dtype", type=str, default="int16", choices=["int16", "int8", "int32"], help="Metadata E datatype")
     parser.add_argument(
         "--bench_torch_sparse",
         type=str,
@@ -253,16 +241,13 @@ def main(
     )
     args = parser.parse_args()
 
-    if args.disable_cache:
-        tilelang.disable_cache()
-
     M, N, K = args.m, args.n, args.k
 
     # Compute total floating-point operations to measure throughput
     total_flops = 2 * M * N * K
 
     # matmul(...) returns (best_latency, best_config, ref_latency)
-    best_result = matmul_sp(M, N, K, T.float16, args.accum_dtype)
+    best_result = matmul_sp(M, N, K, T.float16, args.accum_dtype, e_dtype=args.e_dtype)
     best_latency = best_result.latency
     best_config = best_result.config
     A = torch.randn(M, K, dtype=torch.float16, device="cuda")

benchmark/matmul/benchmark_matmul_sp_compress.py

@@ -0,0 +1,84 @@
+import argparse
+from typing import Optional
+
+import torch
+from tilelang.profiler import do_bench
+from tilelang.utils.sparse import compress, randn_semi_sparse, randint_semi_sparse, torch_compress
+
+SUPPORTED_DTYPE_NAMES = ["float16", "bfloat16", "float32", "int8"]
+SUPPORTED_META_DTYPE_NAMES = ["int8", "int16", "int32"]
+
+
+def _resolve_torch_dtype(name: str) -> torch.dtype:
+    dtype = getattr(torch, name, None)
+    if dtype is None:
+        raise ValueError(f"Unsupported torch dtype: {name}")
+    return dtype
+
+
+def _generate_semi_sparse(m: int, k: int, dtype: torch.dtype, device: str = "cuda") -> torch.Tensor:
+    if dtype in (torch.int8, torch.uint8):
+        return randint_semi_sparse(m, k, low=-64, high=64, dtype=dtype, device=device)
+    return randn_semi_sparse(m, k, dtype=dtype, device=device)
+
+
+def _compress_bytes(input_tensor: torch.Tensor, sparse_tensor: torch.Tensor, meta_tensor: torch.Tensor) -> int:
+    return (
+        input_tensor.numel() * input_tensor.element_size()
+        + sparse_tensor.numel() * sparse_tensor.element_size()
+        + meta_tensor.numel() * meta_tensor.element_size()
+    )
+
+
+def benchmark_compress(
+    m: int,
+    k: int,
+    dtype: torch.dtype,
+    meta_dtype: Optional[torch.dtype] = None,  # noqa: FA100
+):
+    a0 = _generate_semi_sparse(m, k, dtype)
+
+    sparse0, meta0 = compress(a0, meta_dtype=meta_dtype)
+    ref_sparse0, ref_meta0 = torch_compress(a0, meta_dtype=meta_dtype)
+
+    bytes_per_compress = _compress_bytes(a0, sparse0, meta0)
+    bytes_per_torch = _compress_bytes(a0, ref_sparse0, ref_meta0)
+
+    tl_latency_ms = do_bench(lambda: compress(a0, meta_dtype=meta_dtype))
+    torch_latency_ms = do_bench(lambda: torch_compress(a0, meta_dtype=meta_dtype))
+
+    tl_latency_s = tl_latency_ms * 1e-3
+    torch_latency_s = torch_latency_ms * 1e-3
+    tl_throughput_gbps = bytes_per_compress / tl_latency_s / 1e9
+    torch_throughput_gbps = bytes_per_torch / torch_latency_s / 1e9
+
+    return bytes_per_compress, bytes_per_torch, tl_latency_ms, torch_latency_ms, tl_throughput_gbps, torch_throughput_gbps
+
+
+if __name__ == "__main__":
+    parser = argparse.ArgumentParser(description="Benchmark two TileLang compress operators by memory throughput")
+    parser.add_argument("--m", type=int, default=16384, help="Matrix rows")
+    parser.add_argument("--k", type=int, default=16384, help="Matrix columns")
+    parser.add_argument("--dtype", type=str, default="float16", choices=SUPPORTED_DTYPE_NAMES, help="Input dtype")
+    parser.add_argument(
+        "--meta_dtype",
+        type=str,
+        default=None,
+        choices=SUPPORTED_META_DTYPE_NAMES,
+        help="Metadata dtype (defaults to the library choice for the input dtype)",
+    )
+    args = parser.parse_args()
+
+    dtype = _resolve_torch_dtype(args.dtype)
+    meta_dtype = _resolve_torch_dtype(args.meta_dtype) if args.meta_dtype is not None else None
+
+    bytes_per_compress, bytes_per_torch, tl_latency_ms, torch_latency_ms, tl_gbps, torch_gbps = benchmark_compress(
+        args.m,
+        args.k,
+        dtype=dtype,
+        meta_dtype=meta_dtype,
+    )
+
+    print(f"M={args.m} K={args.k} dtype={args.dtype} meta={args.meta_dtype or 'default'}")
+    print(f"tilelang: {tl_latency_ms:.4f} ms, {tl_gbps:.3f} GB/s (bytes={bytes_per_compress})")
+    print(f"torch:    {torch_latency_ms:.4f} ms, {torch_gbps:.3f} GB/s (bytes={bytes_per_torch})")