W8A8 INT8 quantization pipeline

[sandlbn]

Adds end-to-end W8A8 (INT8×INT8) GEMM + fused epilogue dequantization for the transformer attention+FFN pipeline on BMG-G31.

New kernels (include/xe-fuse/kernels/):

gemm_dequant_w8a8.hpp — INT8 GEMM + per-token/per-channel dequant epilogue
gemm_dequant_rope.hpp — above + fused RoPE (K4_W8A8)
gemm_dequant_swiglu.hpp — above + fused SwiGLU/GeGLU (K2_W8A8)
New standalone ops (standalone/ops.hpp, vllm_ops.hpp): dequant_w8a8, dequant_and_rotary_embedding, dequant_and_silu_mul for baseline comparisons.

Activation quantization (compute_rstd.hpp): combined RMSNorm + per-token INT8 quantization in a single 3-pass subgroup kernel.

Pipeline benchmark (autotune/generate_pipeline.py --int8-mode w8a8): three-way comparison — XE_W8A8_FUSED vs VLLM_INT8_EQUIV vs NAIVE_INT8 with per-kernel float reference correctness checks (rtol=0.15).

Results on B70 (LLaMA 3 8B, M=2048): XE_W8A8_FUSED is ~1.9× faster than NAIVE_INT8 and ~1.35× faster than VLLM_INT8_EQUIV.

[Copilot]

Pull request overview

Adds a W8A8 (INT8×INT8) quantized transformer pipeline targeting Xe (BMG-G31), centering on INT8 GEMMs that fuse dequantization and common epilogues (RoPE, SwiGLU/GeGLU), plus baseline kernels and an auto-generated end-to-end benchmark for comparisons.

Changes:

• Introduces fused W8A8 GEMM kernels for dequant-only, dequant+RoPE, and dequant+SwiGLU/GeGLU epilogues.
• Adds standalone INT8 dequant/activation kernels (including “vllm-equivalent” merged kernels) and a fused RMSNorm+INT8 activation quantization kernel.
• Adds a W8A8 pipeline benchmark template + generator option and accompanying tests/docs (incl. Hadamard visitor + correctness test).

Reviewed changes

Copilot reviewed 15 out of 15 changed files in this pull request and generated 5 comments.

Show a summary per file

File	Description
tests/test_k4_w8a8.cpp	New K4 (dequant+RoPE) W8A8 kernel test + reference path
tests/test_k2_w8a8.cpp	New K2 (dequant+SwiGLU) W8A8 kernel test + reference path
tests/test_hadamard_visitor.cpp	New correctness/throughput test for XeHadamardCompute<16>
README.md	Documents new W8A8 pipeline, kernel structs, and baseline ops
include/xe-fuse/visitors/xe_hadamard_compute.hpp	Adds Hadamard epilogue visitor implementation
include/xe-fuse/standalone/vllm_ops.hpp	Adds merged “vllm_int8_equiv” dequant+op kernels
include/xe-fuse/standalone/ops.hpp	Adds naive INT8 activation quantize + dequant baseline ops
include/xe-fuse/kernels/gemm_dequant_w8a8.hpp	Adds fused INT8 GEMM + W8A8 dequant epilogue kernel
include/xe-fuse/kernels/gemm_dequant_swiglu.hpp	Adds fused INT8 GEMM + W8A8 dequant + SwiGLU/GeGLU epilogue kernels
include/xe-fuse/kernels/gemm_dequant_rope.hpp	Adds fused INT8 GEMM + W8A8 dequant + RoPE epilogue kernel
include/xe-fuse/kernels/compute_rstd.hpp	Adds fused RMSNorm + per-token INT8 quantization kernel
include/xe-fuse/builder/epilogue_builder.hpp	Exposes new composed epilogues (DequantRoPE/SwiGLU/GeGLU + HadamardOutput)
autotune/quantize_weights_quaRot.py	Adds offline QuaRot-style weight rotation + INT8 quantization utility
autotune/pipeline_w8a8_template.cpp.j2	Adds auto-generated full W8A8 pipeline benchmark template
autotune/generate_pipeline.py	Adds --int8-mode w8a8 to emit the W8A8 pipeline benchmark

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[gbenms]

code working correctly, some small fixes for issues that may rise later. also, are the hadamard hpp and cpp files used anywhere in the pipeline?

[gbenms]

add: sycl::reqd_sub_group_size(SG_SIZE) or similar, pin the subgroup (SIMD) to SG_SIZE rather than assume the workgroup is a single subgroup. Right now it assumes the 16-wide workgroup is one sub-group, if SIMD8 the reduction only sums half the row.

[gbenms]

add: sycl::reqd_sub_group_size(SG_SIZE) or similar, pin the subgroup (SIMD) to SG_SIZE rather than assume the workgroup is a single subgroup. Right now it assumes the 16-wide workgroup is one sub-group, if SIMD8 the reduction only sums half the row.

[gbenms]

wg_size must be a multiple of the required sub-group size (16):
int wg_size = std::min(256, ((std::max(N / 8, 1) + 15) / 16) * 16)

[gbenms]

wg_size must be a multiple of the required sub-group size (16):
int wg_size = std::min(256, ((std::max(N / 8, 1) + 15) / 16) * 16)

[gbenms]

L>1 isnt treated the same in all places: for buffers L>1 is implemented, and also for token scales, but there isnt really an option to use different weights for L>1 since quantize_weight_cols doenst depent on L, only uses slice 0. no error as long as we use L-1, but would be problematic for larger L. these should be usnified- if the intention is to use the same weights for all batches, they should all read the same weight slice 0, and add here: cute::get<2>(sB_hh) = 0.... should also should fix the token count.

[gbenms]

L>1 isnt treated the same in all places: for buffers L>1 is implemented, and also for token scales, but there isnt really an option to use different weights for L>1 since quantize_weight_cols doenst depent on L, only uses slice 0. no error as long as we use L-1, but would be problematic for larger L. these should be usnified- if the intention is to use the same weights for all batches, the token count should read scale_ch[n], and drop the l*N. also should fix the buffer.