Parent: #204 — Blackwell Phase 2 kernel port
Branch: tairan/blackwell-02-08-int4-fused-silu
Depends on: sub-task 7
Spec: blackwell-kernel-port-v1.md § Sub-task 8
What
Write batchgen_kernels/triton/fused_int4_grouped_silu.py:
V1 approach (split + fused silu_mul):
gate = int4_grouped_gemm(x, Wg)
up = int4_grouped_gemm(x, Wu)
out = silu(gate) * up (via small Triton silu_mul kernel)
Wire into batchgen/moe/fused_int4_wgmma_grouped.py for arch == 'sm100'.
Acceptance
- max_abs_err < 1.0 vs float32 reference
- K2.5 MoE FFN forward completes on B200 without error
Parent: #204 — Blackwell Phase 2 kernel port
Branch: tairan/blackwell-02-08-int4-fused-silu
Depends on: sub-task 7
Spec: blackwell-kernel-port-v1.md § Sub-task 8
What
Write batchgen_kernels/triton/fused_int4_grouped_silu.py:
V1 approach (split + fused silu_mul):
gate = int4_grouped_gemm(x, Wg)
up = int4_grouped_gemm(x, Wu)
out = silu(gate) * up (via small Triton silu_mul kernel)
Wire into batchgen/moe/fused_int4_wgmma_grouped.py for arch == 'sm100'.
Acceptance