[MoE] Add a bf16xfp32 gemm kernel for router gemm#465
Open
jjmiao1 wants to merge 3 commits into
Open
Conversation
Signed-off-by: Avery Miao <avery.miao@intel.com>
Contributor
There was a problem hiding this comment.
Pull request overview
This PR adds a new Intel XPU (Xe2) custom operator, gemm_bf16xfp32, to accelerate MoE router GEMMs where activations are BF16 but router weights remain FP32 for routing stability, by emulating FP32-class accuracy via a dual BF16 GEMM (DualGemm) plus a scaled correction term.
Changes:
- Adds an Xe2 DualGemm kernel (policy-dispatched by M, with split-K + reduction for small-M decode shapes).
- Wires the kernel into the build + PyTorch dispatcher (
torch.ops._xpu_C.gemm_bf16xfp32). - Adds Python helpers, a correctness test against FP32
F.linear, and a benchmark script.
Reviewed changes
Copilot reviewed 14 out of 14 changed files in this pull request and generated 10 comments.
Show a summary per file
| File | Description |
|---|---|
vllm_xpu_kernels/gemm_bf16xfp32.py |
Python API for weight splitting and calling the new op. |
tests/test_gemm_bf16xfp32.py |
Correctness sweep vs FP32 F.linear. |
tests/register_ops.py |
Test helper wrapper for invoking the op via torch.ops._xpu_C. |
benchmark/benchmark_gemm_bf16xfp32.py |
Performance benchmark comparing to FP32 F.linear. |
csrc/xpu/torch_bindings.cpp |
Registers the new op schema + XPU implementation. |
csrc/xpu/gemm_bf16xfp32/gemm_bf16xfp32_interface.h |
Declares the top-level C++ interface. |
csrc/xpu/gemm_bf16xfp32/gemm_bf16xfp32_interface.cpp |
Dispatches to Xe2 implementation (and errors on non-Xe2). |
csrc/xpu/gemm_bf16xfp32/xe_2/CMakeLists.txt |
Adds the Xe2 kernel library target. |
csrc/xpu/gemm_bf16xfp32/xe_2/gemm_bf16xfp32_xe2.h |
Declares the Xe2 entrypoint. |
csrc/xpu/gemm_bf16xfp32/xe_2/gemm_bf16xfp32_xe2.cpp |
Thin wrapper calling the Xe2 implementation. |
csrc/xpu/gemm_bf16xfp32/xe_2/gemm_bf16xfp32_xe2_impl.hpp |
Core Xe2 DualGemm kernel + split-K reduction + policy dispatch. |
csrc/xpu/gemm_bf16xfp32/xe_2/bf16xfp32_epilogue.hpp |
Custom epilogue combining the two accumulator results. |
setup.py |
Installs the additional Xe2 kernel library artifact. |
CMakeLists.txt |
Adds kernel library build + links it into _xpu_C. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Comment on lines
+199
to
+213
| // A [M, K] as (M, Kp, L): row stride = K, L step advances Kp along K. | ||
| StrideA stride_A; | ||
| cute::get<0>(stride_A) = static_cast<int64_t>(K); | ||
| cute::get<2>(stride_A) = static_cast<int64_t>(Kp); | ||
|
|
||
| // B [K, N] as (N, Kp, L): N contiguous, K stride = N, L step = Kp * N. | ||
| StrideB stride_B; | ||
| cute::get<1>(stride_B) = static_cast<int64_t>(N); | ||
| cute::get<2>(stride_B) = static_cast<int64_t>(Kp) * N; | ||
|
|
||
| // D [L, M, N] as (M, N, L): N contiguous, M stride = N, L step = M * N. | ||
| StrideD stride_D; | ||
| cute::get<0>(stride_D) = static_cast<int64_t>(N); | ||
| cute::get<2>(stride_D) = static_cast<int64_t>(M) * N; | ||
|
|
Signed-off-by: Avery Miao <avery.miao@intel.com>
Add Split-K to the mid-M branch, splitting the K reduction over the batch (L) dimension and reducing the [splits, M, N] partials with the existing fused launch_reduce_splits kernel. Mid-M and decode use two constant-free (n_cores-only) split selectors (choose_splits_net / choose_splits_fill); decode keeps its fill-first behavior and prefill stays splits=1. Signed-off-by: Avery Miao <avery.miao@intel.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Purpose
Add an XPU operator
gemm_bf16xfp32that emulates an FP32-precision GEMM (BF16 activation × FP32 weight) using two BF16 GEMM passes (DualGemm). It targets the MoE router GateLinear in models such as Hunyuan-V3 and MiniMax-M3, where the gate weight is kept in FP32 for the numerical stability of routing. Currently this runs as FP32F.linear: the BF16 activation is upcast to FP32 and fed to a oneDNN FP32 matmul. The execution on the FP32 pipe (~12.3 TFLOP/s on B60) is much slower than the BF16 (~98.5 TFLOP/s on B60).gemm_bf16xfp32decomposes the FP32 weight into two BF16 components and combines them in the epilogue, recovering FP32-class accuracy (absolute error ~1e-5) while cutting latency (4–7.7×) on prefill;How it works
The FP32 weight
Wis split into two BF16 tensors:W_high = bf16(W)W_low = bf16((W - fp32(W_high)) / scale)(defaultscale = 1/256)DualGemm The kernel computes
D = X @ W_high + (X @ W_low) * scale, sharing theA(activation) load and fusing the correction in the epilogue. The DualGemm collective is adapted from thedual_gemmexample in sycl-tla: it runs the two BF16 GEMMs from a single kernel soAis loaded once and theW_lowcorrection is applied in the shared epilogue.W_high + W_low * scalereproducesWto ~FP32 precision, so the result can match FP32F.linearto FP32-class accuracy.Split-K for decode At small M (decode) the default tile leaves most Xe-cores idle. For those shapes the kernel splits the K — each slice computes a partial
[M,N], and a small fused kernel sums the partials.Code Changes
csrc/xpu/gemm_bf16xfp32/— kernel interface + implementation (three tile policies: small-M / medium-M / default, dispatched by M; Split-K + fused reduction for decode).vllm_xpu_kernels/gemm_bf16xfp32.py— Python API:split_fp32_weight,gemm_bf16xfp32.tests/test_gemm_bf16xfp32.py— correctness vs FP32F.linear.benchmark/benchmark_gemm_bf16xfp32.py— latency / roofline benchmark vs FP32F.linear.CMakeLists.txt/setup.py— build wiring forgemm_bf16xfp32_xe_2.Usage
Note: Weight preparation is a one-time, offline step.
split_fp32_weightdecomposes the FP32 router weight into the two BF16 components (W_high,W_low); If integrated into upstream vLLM, this will run once at load time in the layer'sprocess_weights_after_loading— the FP32 router weight will split there andW_high/W_lowbe cached.Test Plan
pytest tests/test_gemm_bf16xfp32.py— compares kernel output against FP32F.linearover the M × N × K sweep. Element-wiseassert_close(rtol=1e-3, atol=1e-4)plus aggregate bounds (max_abs < 1e-4,mean_rel < 1e-4) and adual-bf16 ≤ naive-bf16invariant.python benchmark/benchmark_gemm_bf16xfp32.pyon Intel Arc Pro B60 (BMG / Xe2) — reports device + wall latency, TFLOP/s / MFU / GB/s / MBU, and speedup vs FP32.Hardware: Intel Arc Pro B60 (BMG, Xe2). Peaks: BF16 98.5 TFLOP/s, FP32 vector 12.28 TFLOP/s, BW 456 GB/s.
Baseline: FP32
F.linear(oneDNN) — the correctness-equivalent op this kernel replaces.Sweep: M ∈ {2 … 16384}, N ∈ {128, 192, 256}, K ∈ {3072, 4096, 6144}. 100 iters / 20 warmup.
Measurement:
ker dev/fp32 devare on-device kernel time;ker wall/fp32 wallare end-to-end eager wall-clock;dev speedup/wall speeduparefp32 / kernel. TFLOP/s counts both BF16 passes (4·M·N·K); MFU vs BF16 peak; GB/s / MBU from IO traffic (A + W_high + W_low + D) vs BW peak.Test Result
pytest tests/test_gemm_bf16xfp32.py: all passAccuracy is measured against the FP32 (oneDNN) baseline across the full sweep: absolute error ~1e-5 (worst 2.5e-5), mean relative ~2e-5.
F.linear(see tables below)Representative MNK— Hunyuan-V3 router gemm shape (N = 192 experts, K = 4096 hidden), full M sweep:
Representative MNK — MiniMax-M3 router shape (N = 128 experts, K = 6144 hidden), full M sweep:
Benchmark Results Summary
F.linearis stuck on the FP32 vector pipe (12.28 TFLOP/s). Result: 4–7.7× device speedup, reaching up to ~89% MFU.