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

Activity

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

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