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
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:
Wire into batchgen/moe/fused_int4_wgmma_grouped.py for arch == 'sm100'.
Acceptance