Skip to content

[blackwell] sm100: INT4 grouped GEMM Triton kernel (K2.5 decode) #230

Description

@Andrewxu313

Parent: #204 — Blackwell Phase 2 kernel port
Branch: tairan/blackwell-02-07-int4-grouped-gemm
Spec: blackwell-kernel-port-v1.md § Sub-task 7

What

Write batchgen_kernels/triton/int4_grouped_gemm.py:

  • INT4 x BF16 grouped GEMM for Kimi K2.5 MoE decode
  • Inline dequant: (w_int4 - 8).to(bf16) * scale_per_group, group_size=128
  • Small-M regime (M <= 32 per expert in decode)
  • autotune with BLOCK_M=16/32 configs

Wire into batchgen/moe/fused_int4_wgmma_grouped.py for arch == 'sm100'.

Acceptance

  • max_abs_err < 0.5 vs float reference (INT4 quantization noise expected)
  • Throughput >= 50% of Marlin at decode shapes

Metadata

Metadata

Assignees

No one assigned

    Labels

    enhancementNew feature or request

    Type

    No type

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions